cuda by example第10章 流 学习笔记

由于第10章代码相对简单，如果前几章代码没有问题的话这章代码问题应该不大，但这章主要涉及了并行的概念，因此这篇博客主要谈一下自己对书中并行操作的理解，如有不对的地方请多指正

1.cudaHostAlloc（）

这个函数与malloc（）的主要区别就是用cudaHostAlloc（）分配的内存将被锁定，他不会被分页并且与硬盘中其他页进行交换，这里涉及到一点操作系统的知识，我们知道内存的空间大小是有限的，而硬盘中的数据远大于内存的大小，因此我们通常会使用虚拟内存技术将cpu，gpu当前需要的数据放入内存中，将暂时不用的数据放在硬盘当中，当需要用到硬盘中的数据时再将内存中的数据换出。这样好像是扩充了内存的大小但并没有实际扩大内存大小，因此被称为虚拟内存技术。而每次交换数据是是以页为单位进行交换。因此可以将页理解为一种单位大小，类似一块数据。

再解释完上面的概念以后我们来理解一下为什么要使用cudaHostAlloc（）分配内存，因为GPU可以直接通过DMA技术获取到内存中的数据而在传输过程中不需要cpu的介入，因此当GPU通过DMA技术获取数据时，cpu将这块内存将这块内存中的数据移动到硬盘上，这样的话DMA会终止传输数据，向cpu发出请求将这块数据重新从硬盘移动到内存中，这样就增加了数据传输所需要的时间，而在使用cudaHostAlloc（）分配内存后，这块内存不会被替换掉，变相的加快了速度。

书中分别给了使用malloc与cudaHostAlloc两种方法传输数据的性能发现使用cudaHostAlloc有显著的性能提升。

流

流可以理解为是一个执行队列，用于在设备上并行处理任务。所有提交到同一流的操作都会按顺序执行，而不同流中的操作可以并行进行。

单个cuda流

书中举了一个计算a数组，b数组相加的平均值并存入数组c中的例子。由于我们这次使用的数据比较大，没办法在GPU中一次性输入所有数据，因此我们需要把要输入的数据分成一个个小的块，进行分块计算。

每块计算可以分为3个步骤：

1.将这一块数据复制到GPU上

2.在GPU上执行kernel函数

3.将计算的结果返回主机当中

代码

for (int i=0; i<FULL_DATA_SIZE; i+= N) {
        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream ) );

        kernel<<<N/256,256,0,stream>>>( dev_a, dev_b, dev_c );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream ) );

    }

我们发现这段代码与之前的不同之处主要是使用了cudaMemcpyAsync（）这个函数，注意这个函数的最后一个参数我们使用了一个stream流，这个流相当于一个队列，我们把第一个操作（host_a+i这个内存中的数据放入dev_a中）放入队列中，再将第二个操作（host_b+i这个内存中的数据放入dev_b中）放入队列中，后执行核函数（第三步操作，发现核函数中也有一个stream参数，相当于把这个操作也放入流中）再将第四个操作（将计算得到的存储在gpu中的数据返回到主机中）放入到队列中。流就像一个有序的工作队列，gpu依次执行内部的操作

 HANDLE_ERROR( cudaStreamSynchronize( stream ) );

一个同步操作确保stream流中的操作全部完成

多cuda流

书中关于多cuda流的使用给了两个例子，一种是改进前，一种是改进后，我们这里先讨论该进前的例子

改进前

我们发现这个图中给出了两个流的执行时间图

以下是执行这个时间图的代码

  for (int i=0; i<FULL_DATA_SIZE; i+= N*2) {
        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a0, host_a+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b0, host_b+i,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream0 ) );

        kernel<<<N/256,256,0,stream0>>>( dev_a0, dev_b0, dev_c0 );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i, dev_c0,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream0 ) );


        // copy the locked memory to the device, async
        HANDLE_ERROR( cudaMemcpyAsync( dev_a1, host_a+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );
        HANDLE_ERROR( cudaMemcpyAsync( dev_b1, host_b+i+N,
                                       N * sizeof(int),
                                       cudaMemcpyHostToDevice,
                                       stream1 ) );

        kernel<<<N/256,256,0,stream1>>>( dev_a1, dev_b1, dev_c1 );

        // copy the data from device to locked memory
        HANDLE_ERROR( cudaMemcpyAsync( host_c+i+N, dev_c1,
                                       N * sizeof(int),
                                       cudaMemcpyDeviceToHost,
                                       stream1 ) );
    }

但是我们执行之后发现所耗时间并没有明显减小，这引出了我们接下来要介绍的GPU的工作原理

GPU工作调度机制

GPU引擎实际上分为一个内存复制引擎还有一个核函数执行引擎下图为改进前我们的操作对应的gpu操作

对两个引擎来说他们执行的操作以及顺序如下图所示

我们发现改进前的方法在执行第0个流的复制c操作时必须等待第0个流执行核函数后才能执行，这也解释了为什么时间并没有多少的提升

改进后

改进的方法就是尽可能加大“宽度”，可以理解为尽可能多的让stream0与stream1尽可能多的并行操作以下是书中所给的示例图

我们发现在第1个流执行复制c操作时只需要等待第1个流执行核函数后，并不需要等待第0个流所有的操作执行完成后这样减少了所消耗的时间。

结尾

通过对流的理解我们应该对cuda并行操作有了一定的理解，我们发现GPU引擎实际上被分为了两个部分，不同的操作需要使用到不同的引擎因此，合理的使用引擎增加引擎并用的时间是提升效率的关键。以上仅代表个人理解，如果有什么不对的地方请评论区指正😘