CUDA Stream相关知识

流的分类

学习知识一般有两个方向，一个是往底层钻研，达到庖丁解牛的程度，还有一个方向是往上钻研，达到高屋建瓴的本事，二者缺一不可。今天这个流概念就是稍微上层的思想。流可以理解为一个管道，码农所处在的位置就是CPU, 现在你想往GPU发任务，那么你就要先建立一个管道，然后往这个管道塞入你想做的事情，GPU然后会从管道的另一头获取任务，所以同一个管道中的API严格按照顺序执行。

流可以分为以下两个类型:(其实在我看来就是普通的流，只不过是如果我们自己不创建的话，系统会给我们默认生成一个流，这个流叫做空流，不过这个空流有个特性就是会阻塞其他的流，后面会具体解释。)

• 空流
• 非空流
  我们一般写的cuda代码都是用的空流，也就是程序生成的一个默认流。下面是一个最简单的流中的操作，伪代码如下:

memcpy(dst,src,size,HostToDevice);
kernel<<<M,N>>>(arg_1; arg_2);
memcpy(dst,src,size,DeviceToHost);

根据执行的时序图可以看到，memcpy会是host端阻塞，而kernel被调用后会立马控制权会立马发回给host, 而接下来的内存拷贝又会阻塞进程，以上就是一个初级的流原理。

流的调度

从概念上来看，两个不同的流中的kernel是可以在物理硬件上并行的，但是实际往往会复杂的多。

第一阶段

一开始的Fermi GPU是支持16个kernel同时进行的，但是底层硬件只有一个工作队列，所以到了底层执行的时候几乎还是串行的，比如有两个流1，2，其中1：a->b->c, 2流的kernel:d->e->f, 当进入底层硬件时候，还是a->b->c d–>e->f, 比如在执行b的时候，硬件会看看b的依赖a是不是执行完了，执行完了才执行b, b执行完了后，c发现他的依赖都执行完了，然后就会执行，而d发现他的依赖也执行完了（因为他压根没有依赖），所以真正并行的只有c和d两个依赖。这个问题的本质还是由于硬件只有一个工作队列导致的，所以所谓的kernel并行不过就是一句空话。

第二阶段

第二个阶段出现了Hyper-Q技术，这个技术的核心就是出现了多个硬件工作队列, 这个就可以让多个cpu线程，或者进程在同一个GPU上搞事情了。
kerpler有32个硬件工作队列，如果创建的流超过32个的话，会有部分流共享一个硬件工作队列。