CUDA流任务并行

       CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。可以将每个流视为GPU的一个任务，并且这些任务可以并行执行，即相同流顺序执行，不同流并行执行；不同流并行执行时不同流所要执行的任务要没有依赖关系；当不手动创建流时，cuda将会默认一个流操作。

       在硬件选择上，这里有一个概念，支持设备重叠功能，支持设备重叠功能的 GPU 能够在执行一个 CUDA C 核函数的同时，还能在设备和主机之间执行复制操作；这在流并行过程中很重要，我们假设有流A和流B，设备重叠就会允许流A在复制过程中同时流B进行核函数计算，这会大大加快速度；

cudaDeviceProp   prop;
int  whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDevice(&prop, whichDevice);
if(prop.deviceOverlap){
   std::cout<<"the device  will handle overlaps"<<std::endl;
}

        我们知道cudaMemcpy与CPU操作是同步的，为了实现设备重叠，cuda提供了cudaMemcpyAsync用于数据拷贝操作,它是异步的，不会等待复制完成就会执行程序的下一步；

        但注意，cudaMemcpyAsync仅对分页锁定的主存储器有效，如果传入指向可分页存储器的指针，那么将返回一个错误；

        页锁定的主机内存由cudaHostAlloc()分配。页锁定的主机内存也称为固定内存或不可分页内存，它的重要属性就是：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全的使用应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。事实上，当使用可分页内存进行复制时，复制操作将执行两遍，第一遍从可分页内存复制到一块“临时的”页锁定内存，然后再从这个页锁定内存复制到GPU上。因此，当在GPU和主机间复制数据时，这种差异会使也锁定主机内存的性能比标准可分页内存的性能要高大约2倍。然而，我们也不能进入另一个极端：查找每一个malloc调用并将其替换为cudaHostAlloc调用。固定内存是一把双刃剑，当使用固定内存是，你将失去虚拟内存的所用功能。特别是，应用程序中使用每个页锁定内存时都需要分配物理内存，因为这些内存不能交换到磁盘上。这意味着，与使用标准的malloc调用相比，系统将更快的耗尽内存（概念选自《GPU高性能编程CUDA实战》）。

        从上，页锁定内存不仅在主机与设备之间复制数据快，而且在流并行中扮演着重要的作用；

多个CUDA流宽度优先而非深度优先      

     深度优先就是程序按顺序把一个流的操作添加之后再添加下一个流操作,如下：

for(int i =0; i< FULL_DATA_SIZE; i+= 2*N){
    cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0);
    cudamemcpyAsync(dev_b0, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0);
    kernel<<<N/256,256,0,stream0>>>(dev_a0, dev_b0,dev_c0);
    cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int),cudaMemcpyDeviceToHost, stream0);
    cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1);
    cudamemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1);
    kernel<<<N/256,256,0,stream1>>>(dev_a1, dev_b1,dev_c1);
    cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int),cudaMemcpyDeviceToHost, stream1);
}

先添加stream0再添加stream1操作,按流模型来说，因为这里的拷贝操作还是核函数都是异步的，stream0和stream1并行计算，速度应该比用单个流提了不少，但是却不然，为什么呢？这就涉及到GPU的硬件调度：

 在硬件中并没有流的概念，而是包含一个或多个引擎（主机到设备，设备到主机可能是分开的两个引擎）来执行内存复制操作，以及一个引擎来执行核函数。这些引擎彼此独立的对操作进行排队；

也就是说内存复制和核函数在GPU上是不同的引擎在执行，那么在同一流上核函数和复制操作相邻时，就会发生一个现象：两个操作在不同的引擎上，但是流模型又要保证同一流上两个程序执行的先后顺序，那么怎么办？cuda驱动程序为了保证硬件的执行单元不破坏流之间的依赖性，将前一个操作阻塞，等待完成后，再继续执行后一个；

所以为了高效利用CUDA流，提出了流宽度优先的概念，即将两个流之间的操作交叉添加：

for(int i =0; i< FULL_DATA_SIZE; i+= 2*N){
    cudaMemcpyAsync(dev_a0, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0);
    cudaMemcpyAsync(dev_a1, host_a + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1);
    cudamemcpyAsync(dev_b0, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream0);
    cudamemcpyAsync(dev_b1, host_b + i + N, N*sizeof(int),cudaMemcpyHostToDevice, stream1);
    kernel<<<N/256,256,0,stream0>>>(dev_a0, dev_b0,dev_c0);
    kernel<<<N/256,256,0,stream1>>>(dev_a1, dev_b1,dev_c1);
    cudaMemcpyAsync(host_c + i, dev_c0, N*sizeof(int),cudaMemcpyDeviceToHost, stream0);
    cudaMemcpyAsync(host_c + i + N, dev_c1, N*sizeof(int),cudaMemcpyDeviceToHost, stream1);
}

   http://www.mamicode.com/info-detail-1770665.html