TensorFlow和pytorch中的pin_memory和non_blocking设置是做什么的,又是否有用???(续)
接上篇:
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
还是
的区别;如果non_blocking=True,所包装的CUDA执行的语句就是cudaMemcpyAsync
,如果non_blocking=False,所包装的CUDA执行的语句就是cudaMemcpyAsync
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解放出来,进一步提高效率,可以使用
指令,这样CPU端也不需要等待GPU端对这条语句执行完就可以继续执行,而GPU端执行这条语句的同时其他stream队列中的操作也可以同时并行(需要注意的是在GPU的多个stream队列中同一时刻只能有一个cudaMemcpyAsync
指令在执行,如果其他stream队列中也有cudaMemcpyAsync
指令需要执行那么则需要等待这条cudaMemcpyAsync
指令执行完才可以执行,单方向的cudaMemcpyAsync
一个时刻只能有一个,但是双方向copy的可以同时运行cudaMemcpyAsync
)。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端的操作顺序。
虽然使用
可以实现CPU和GPU的同时运行,但是如果CPU中之后的运行中需要使用GPU通过cudaMemcpyAsync
拷贝回host内存中的数据,那么我们就无法知道此时的数据是否是GPU中执行完操作copy回来的,为此我们就需要手动使用同步操作,保证内存中的数据是GPU拷贝回的数据,为此我们可以保证CPU和GPU同步的操作有:cudaMemcpyAsync
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队列。
====================================================
posted on 2022-11-12 20:10 Angry_Panda 阅读(313) 评论(0) 编辑 收藏 举报