创建流
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的核函数。