CUDA： 流

1. 页锁定主机内存

　　c库函数malloc()分配标准的，可分页(Pagable)的内存，cudaHostAlloc()分配页锁定的主机内存。页锁定内存也称为固定内存（Pinned Memory）或者不可分页内存，它有个重要属性：操作系统将不会对这块内存分页并交换到磁盘上，从而确保了该内存始终驻留在物理内存中。因此，操作系统能够安全的使某个应用程序访问该内存的物理地址，因为这块内存将不会被破坏或者重新定位。

　　由于GPU知道内存的物理地址，因此可以通过“直接内存访问（Direct Memory Access ,DMA）”技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU的介入，这也就同样意味着，CPU可能在DMA的执行过程中将目标内存交换到磁盘上，或者通过更新操作系统的可分页表来重新定位目标内存的物理地址。CPU可能会移动可分页的数据，这就可能对DMA操作造成延迟。因此，在DMA复制过程中使用固定内存是非常重要的。事实上，当使用可分页内存进行复制时，CUDA驱动程序仍然会通过DMA把数据传输给GPU。因此，复制操作将执行两遍，第一遍从可分页内存复制到一块临时的页锁定内存，然后再从这个页锁定内存复制到GPU上。因此，每当从可分页内存中执行复制操作时，复制速度将受限于PCIE传输速度和系统前端总线速度相对较低的一方。

　　当使用固定内存时，将失去虚拟内存的所有功能，导致系统更快的耗尽内存，降低系统整体性能。建议，仅对cudaMemcpy（）调用中的源内存或者目标内存，才使用锁定内存，并且不再需要使用它们时立即释放，而不是等到应用程序关闭时才释放。

2. CUDA 流

　　CUDA流表示一个GPU操作队列，并且该队列中的操作将以指定的顺序执行。可以将每个流视为GPU的一个任务，并且这些任务可以并行执行。

选择一个支持设备重叠功能的设备:

cudaDeviceProp   prop;
int    whichDevice;
cudaGetDevice(&whichDevice);
cudaGetDevice(&prop, whichDevice);
if(!prop.deviceOverlap){
  printf("Device will not handle overlaps , so no speed up from streams
");  
}

创建流：

cudaStream_t     stream;
cudaStreamCreate(&stream);

使用流：

for(int i =0; i< FULL_DATA_SIZE; i+=N){
    cudaMemcpyAsync(dev_a, host_a + i, N*sizeof(int),cudaMemcpyHostToDevice, stream);
    cudamemcpyAsync(dev_b, host_b + i, N*sizeof(int),cudaMemcpyHostToDevice, stream);
    kernel<<<N/256,256,0,stream>>>(dev_a, dev_b,dev_c);
    cudaMemcpyAsync(host_c + i, dev_c, N*sizeof(int),cudaMemcpyDeviceToHost, stream);   
}

任何传递给cudaMemcpyAsync的主机内存指针都必须已经通过cudaHostAlloc分配好内存，也就是说，你只能以异步方式对页锁定内存进行复制操作。

GPU与主机同步,调用cudaStreamSynchronize()并制定想要等待的流：

cudaStreamSynchronize(stream);

释放流：

cudaStreamDestroy(stream)

3. 使用多个流

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);
}

cudaStreamSynchronize(stream0);
cudaStreamSynchronize(stream1);

4.GPU的工作调度机制

　　在硬件中并没有流的概念，而是包含一个或多个引擎（主机到设备，设备到主机可能是分开的两个引擎）来执行内存复制操作，以及一个引擎来执行核函数。这些引擎彼此独立的对操作进行排队，因此导致下图所示任务调度情形：

应用程序首先将第0个流的所有操作放入队列，然后是第一个流的所有操作。CUDA驱动程序负责按照这些操作的顺序把他们调度到硬件上执行，这就维持了流内部的依赖性。图10.3说明了这些依赖性，箭头表示复制操作要等核函数执行完成之后才能开始。

于是得到这些操作在硬件上执行的时间线：

由于第0个流中将c复制回主机的操作要等待核函数执行完成，因此第1个流中将a和b复制到GPU的操作虽然是完全独立的，但却被阻塞了，这是因为GPU引擎是按照指定的顺序来执行工作。记住，硬件在处理内存复制和核函数执行时分别采用了不同的引擎，因此我们需要知道，将操作放入流队列中的顺序将影响着CUDA驱动程序调度这些操作以及执行的方式。

5. 高效使用多个流

如果同时调度某个流的所有操作，那么很容易在无意中阻塞另一个流的复制操作或者核函数执行。要解决这个问题，在将操作放入流的队列时应采用宽度优先方式，而非深度优先方式。如下代码所示：

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);
}

如果内存复制操作的时间与核函数执行的时间大致相当，那么新的执行时间线将如图10.5所示，在新的调度顺序中，依赖性仍然能得到满足：

由于采用了宽度优先方式将操作放入各个流的队列中，因此第0个流对c的复制操作将不会阻塞第1个流对a和b的内存复制操作。这使得GPU能够并行的执行复制操作和核函数，从而使应用程序的运行速度显著加快。