TensorFlow和pytorch中的pin_memory和non_blocking设置是做什么的，又是否有用？？？（续）Tenso

接上篇：

​​TensorFlow和pytorch中的pin_memory和non_blocking设置是做什么的，又是否有用？？？​​

参考：

​​How to Overlap Data Transfers in CUDA C/C++​​

如何实现nvidia显卡的cuda的多kernel并发执行？？？

=================================================

首先要说一下，CPU的运行原理大家基本都有一个大致的了解，对于计算机科班出身的人来说操作系统也是必修的主修课程，但是GPU的运行原理却很少有人知道，而且GPU这个东西本身有很多的技术都是不公开的，尤其是GPU做通用计算的基本上都是那几个公司在搞，尤其是对于软件工程师，哪怕是CUDA领域的软件工程师，对于GPU的很多内部运行原理也是不清除的，对于普通的CUDA用户来说基本可以说只需要了解CUDA runtime api和driver api这两个层面的接口即可，可以说CUDA的运行原理对于软件人员来说基本是透明的；但是GPU编程不同于CPU编程，编码人员还需要对GPU的运算单元和内存结构层次有个了解才行，不然SIMP（单指令多线程）的代码是搞不好的。

接上文，本篇主要是要说下深度学习框架中的non_blocking设置是做什么的，又是否有用？

对这个non_blocking要有所了解主要还是要知道两个概念，一个是pinned memory；一个是stream队列。关于pinned memory上篇post已经给了介绍，这里再概要的复述下：host向device传输数据时（64kb大小以上数据），需要CPU在内存中指定一块pinned memory，将需要传输的数据放入到这块pinned memory中，然后向GPU发送指令要其将这块pinned memory中的数据copy回GPU的显存中。（关于pinned memory的介绍见：​​TensorFlow和pytorch中的pin_memory和non_blocking设置是做什么的，又是否有用？？？​​）

下面说下stream队列。首先我们需要知道一点，那就是CPU中运行的代码和数据都是从硬盘上调入到内存中的，不同进程被CPU调度执行都是需要从内存中读取对应的运行代码和数据的，而GPU是不能从硬盘中读取运行代码和数据的，这也是为啥GPU端叫做device，CPU端叫做host的原因。可以说GPU中的运行代码和数据都是从host端内存导入的，GPU中每一步操作指令都是CPU端从host内存中发送给GPU的（个人认为发送给GPU的函数运行指令只需要发送一次，然后就可以保存到GPU端的显存空间中，维护一个CPU进程对应在GPU端的显存空间都是被其对应的context所维护的）。为了理解上更加简单，我们可以暂且认为GPU端需要运行的指令都是CPU端实时发送来的，GPU端将收到的运行指令存放在队列中，这样的队列就叫做stream队列。

CPU端向GPU发送运行指令时，如果没有特殊的指定，所有的运行执行都是被GPU存放到显存中该CPU进程所对应的context下的default stream队列中（为了理解方便，我们假设每个CPU进程调用CUDA时在GPU上都只开一个context），当然如果在CPU向GPU发送指令时明确的标明该指令需要存放在某个非default的stream队列中也是可以的。GPU具体运行时是需要从stream队列中取出具体的运行指令的，由于GPU是有着成百上千个运算核心的，因此GPU可以同时运行多个stream队列中的运行指令的，但是这里需要注意几点：

1. default stream队列独占整个context（也是独占整个GPU），也就是说non-default stream队列中个有指令在运行的时候default stream中的指令是不会执行的，而default stream中的指令执行时non-default stream中的指令是不能执行的，也就是说default stream中的指令会被其他non-default stream中的指令所阻塞，同样default stream中的指令也会堵塞其他non-default stream中的指令；

2. 如果CPU端发送给GPU的指令为kernel类别的话，CPU进程发送完指令即可继续执行不需要等待GPU端对该条指令的执行完，也就是说CPU端的执行和GPU端kernel的执行是可以并行的。

3. 如果CPU端发送给GPU端的指令为​​cudaMemcpy​​，如果待传输的数据是pageable memory，那么CPU端需要在host内存中临时申请pinned memory，把待传输的数据copy到这个pinned memory，然后发送指令给GPU，要GPU将这块pinned memory中的数据copy走，然后CPU端陷入阻塞状态，等待GPU端copy任务结束再唤醒，而GPU端收到这条copy指令后把该指令存入到default stream队列中，由于GPU端的指令执行是需要排队的，所以CPU需要一直等待GPU端执行该指令并完成设备间数据copy后才可以继续执行。

4. 如果CPU端发送给GPU端的指令为cudaMemcpyAsync，与cudaMemcpy指令相同的是如果待传输的数据是pageable memory，那么CPU端需>要在host内存中临时申请pinned memory，把待传输的数据copy到这个pinned memory，然后发送指令给GPU，要GPU将这块pinned memory中的数据copy走，但是与cudaMemcpy所不同的是CPU端发送给GPU端执行后CPU端不陷入阻塞状态继续向下执行；而且cudaMemcpyAsync指>令与cudaMemcpy指令不同的是这条指令可以被指定进入GPU端的哪个stream队列中，这样GPU端执行这条指令时就不会阻塞GPU端其他non-default stream队列中的指令执行了。

5. default stream队列的指令不能和其他non-default stream队列中的指令同时执行，多个non-default stream中的指令可以同时执行。

可以说，在深度学习框架中是否使用non_blocking的区别其本质就是在host与device间copy数据时CUDA执行是使用​​cudaMemcpy​​还是​​cudaMemcpyAsync​​的区别；如果non_blocking=True，所包装的CUDA执行的语句就是​​cudaMemcpyAsync​​，如果non_blocking=False，所包装的CUDA执行的语句就是​​cudaMemcpy​​。

===================================

How to Overlap Data Transfers in CUDA C/C++中的一些例子：

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);
myCpuFunction(b)

第一条语句执行，CPU发送指令给GPU后会阻塞直到GPU端的操作结束；第二条语句执行，CPU发送给GPU命令后直接向下执行第三条语句；第三条语句执行，CPU发送指令给GPU后会阻塞直到GPU端的操作结束；第四条语句执行，CPU执行操作不和GPU有任何关系。

需要注意的是cudaMemcpy这个操作的执行，首先CPU端发送这条执行指令给GPU后会阻塞直到GPU端执行完这条指令；同时，GPU端排队到执行这条default stream中的语句是会阻塞其他stream 队列中的指令执行；也就是说GPU端执行这个语句时会把之前所有发送到GPU端的指令全部执行完，然后才会执行这条语句，而CPU端只有等到GPU上这条语句执行结束才可以继续向下执行，这就意味着cudaMemcpy这个指令的执行implicit（隐式）的将CPU端和GPU端进行了同步，可以说cudaMemcpy是非常简单粗暴的实现了CPU端和GPU端的同步操作。

从上面的代码中可以看到第二条语句和第一条语句有依赖关系，第三条语句与第二条语句有依赖关系，但是第四条语句和前三天语句均没有依赖关系；同时由于第二条语句的执行并不会阻塞CPU端的执行，而第三条语句的执行会阻塞CPU的执行，因此完全可以把第四条语句发到第三条语句前执行，这样就实现了CPU端和GPU端同时运行，提高效率，修改后的代码如下：

cudaMemcpy(d_a, a, numBytes, cudaMemcpyHostToDevice);
increment<<<1,N>>>(d_a)
myCpuFunction(b)
cudaMemcpy(a, d_a, numBytes, cudaMemcpyDeviceToHost);

上面代码中CPU发送increment指令给GPU后就可以执行myCpuFunction操作，这样CPU端执行myCpuFunction函数的同时GPU端也可以同时执行increment函数，这样就实现了CPU和GPU的并行，提高运行效率。

为了将CPU解放出来，进一步提高效率，可以使用​​cudaMemcpyAsync​​指令，这样CPU端也不需要等待GPU端对这条语句执行完就可以继续执行，而GPU端执行这条语句的同时其他stream队列中的操作也可以同时并行（需要注意的是在GPU的多个stream队列中同一时刻只能有一个​​cudaMemcpyAsync​​指令在执行，如果其他stream队列中也有​​cudaMemcpyAsync​​指令需要执行那么则需要等待这条​​cudaMemcpyAsync​​指令执行完才可以执行，单方向的​​cudaMemcpyAsync​​一个时刻只能有一个，但是双方向copy的可以同时运行​​cudaMemcpyAsync​​）。使用​​cudaMemcpyAsync​​，上面代码可以写为：

cudaStream_t stream1;
cudaError_t result;
result = cudaStreamCreate(&stream1);
result = cudaMemcpyAsync(d_a, a, N, cudaMemcpyHostToDevice, stream1);
increment<<<1,N,0,stream1>>>(d_a);
result = cudaMemcpyAsync(a, d_a, N, cudaMemcpyDeviceToHost, stream1);
result = cudaStreamDestroy(stream1);

myCpuFunction(b)

上面代码中CPU端和GPU端可以实现并行，同时由于GPU端的所有指令都放在了相同的non-default stream队列中，因此保证了GPU端的操作顺序。

虽然使用​​cudaMemcpyAsync​​可以实现CPU和GPU的同时运行，但是如果CPU中之后的运行中需要使用GPU通过​​cudaMemcpyAsync​​拷贝回host内存中的数据，那么我们就无法知道此时的数据是否是GPU中执行完操作copy回来的，为此我们就需要手动使用同步操作，保证内存中的数据是GPU拷贝回的数据，为此我们可以保证CPU和GPU同步的操作有：

cudaDeviceSynchronize()  阻塞CPU直到之前发送到GPU的所有指令执行结束，该操作的同步功能与cudaMemcpy相同；

cudaStreamSynchronize(stream)  阻塞CPU直到之前发送到GPU上指定的stream队列中的所有指令执行结束，该操作不影响GPU上其他stream队列中的指令执行；

cudaEventSynchronize(event)    event在record的时候是需要指定stream队列的，因此该语句执行时阻塞CPU直到之前发送到GPU上指>定的stream队列中的record event的指令执行结束；

相关的：

cudaStreamWaitEvent(event)               与上面三个阻塞host端CPU的不同，该语句阻塞的是某个stream队列，也就是说CPU执行该语句是将执行发送给某个stream队列，该stream队列执行这个执行后便会阻塞，直到这个GPU上的event事件被record，而这个event是其他stream队列中的，而且这个其他stream队列也可以是其他GPU上的stream队列。

====================================================