cuda流硬件原理和多流优化

创建流

    cudaStream_t stream;
    cudaStreamCreate(&stream);

kernel<<<>>>分配符第四个参数可带stream，GPU将顺序执行已经放到stream中的工作。

fun<<<16, 256, 0, stream>>>(d1_data, d2_data);

cudaMemcpyAsync以异步方式进行拷贝，将工作交到stream中。

cudaMemcpyAsync(h1_data, d2_data, PRO_MAX_DATA, cudaMemcpyDeviceToDevice, stream);

cudaStreamSynchronize阻塞等待流执行完成。

cudaStreamSynchronize(stream);

一些新的NVIDIA GPU支持同时一个核函数和两次内存复制操作，一次是从主机到设备，另一次是从设备到主机。在任何支持内存复制和核函数的执行相互重叠的设备.上，当使用多个流时，应用程序的整体性能都会提升。

假设工作为将主机上的a数组+b数组放到c数组。由于GPU支持同时进行核函数和两种复制，所以可以用两个流做处理。主要代码中，for遍历总数据，在for中：

一、将a部分数据复制到GPU，
二、将b部分数据复制到GPU，
三、调用内核函数计算，
四、将结果复制到c。
上述四个步骤依次放到两个流中，保证所有任务顺序执行。

并且当GPU在执行四的时候，CPU已经开始执行下一次的步骤一二三了，使流中一直有任务，提高利用率。

但由于GPU在硬件上，一个引擎负责复制任务，一个引擎负责核函数任务，所有的复制和核函数分别在两个引擎上排队。并且步骤四对步骤三有依赖性，所以流1的步骤四会等待步骤三结束，流2的复制任务在同一引擎排在步骤四后面，也必须等待流1的步骤四结束，所以阻塞。

更好的方法是先将两个流的复制到设备放到复制引擎，最后进行核函数和复制回主机，这样就能在流1执行完核函数，开始步骤四时，核函数引擎同时开始执行流2的核函数。