3.2.6.5 流
应用通过流管理并发。流是一系列顺序执行的命令。另外，流之间相对无序的或并发的执行它们的命令;这种行为是没有保证的，而且不能作为正确性的的保证(如内核间的通信没有定义)。
可以通过创建流对象来定义流，且可指定它作为一系列内核发射和设备主机间存储器拷贝的流参数。下面的代码创建了两个流且在分页锁定存储器中分配了一个名为hostPtr的浮点数组。
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
下面的代码定义的每个流是一个由一次主机到设备的传输，一次内核发射，一次设备到主机的传输组成的系列。
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size,
size, cudaMemcpyHostToDevice, stream[i]);
for (int i = 0; i < 2; ++i)
MyKernel<<<100, 512, 0, stream[i]>>> (outputDevPtr + i * size, inputDevPtr + i * size, size);
for (int i = 0; i < 2; ++i)
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost,
每个流将它的hostPtr输入数组的部分拷贝到设备存储器数组inputdevPtr，调用MyKernel()内核处理inputDevPtr，然后将结果outputDevPtr传输回hostPtr同样的部分。使用两个流处理hostPtr允许一个流的传输和另一个流的执行重叠。为了使用重叠hostPtr必须指向分页锁定主机存储器。
最后调用的cudaThreadSynchronize()保证所有的流在进一步执行前已经完成。这个函数强制运行时等待所有流中的任务都完成。cudaStreamSynchronize()强制运行时等待某个流中的任务都完成。可用于同步主机和特定流，同时允许其它流继续执行。cudaStreamQuery()用于查询流中的所有之前的命令是否已经完成。为了避免不必要的性能损失，这些函数最好用于计时或隔离失败的发射或存储器拷贝。
调用cudaStreamDestroy()来释放流。
cudaStreamDestroy(stream[i]);
cudaStreamDestroy()等待指定流中所有在前的任务完成，然后释放流并将控制权返回给主机线程。
如果是下面情况，来自不同流的两个命令也不能并发，或分页锁定主机存储器分配，设备存储器分配，设备存储器设置，设备之间拷贝，或主机线程在0号流中调用的任何它们之间的CUDA命令(包含没有指定任何流参数的内核发射和设备和主机间存储器拷贝)。
任何需要依赖检测以确定内核发射是否完成的操作会阻塞CUDA上下文中后面任何流中所有的内核发射直至被检测的内核发射完成。需要依赖检测的操作包括同一个流中的一些其它类似被检查的发射的命令和流中的任何cudaStreamQuery()调用。因此，应用应当遵守这些指导以提升潜在的内核并发执行：
1、所有独立操作应当在依赖操作之前发出。
2、任何类型同步尽量延后。
G.4.1节描述的一级缓存和共享存储器的配置切换为所有未完成的内核发射插入了一个设备端同步栅栏。
