本文从CSDN上转移过来：
http://blog.csdn.net/mounty_fsc/article/details/51092933
1.流（stream）的理解
一个流对应并发的概念，是一组顺序执行的操作（可能由多个主机线程发出）；
多个流对应并行的概念，因为发生顺序具有不确定性。
2.相关函数
//基本函数
cudaStream_t stream//定义流
cudaStreamCreate(cudaStream_t * s)//创建流
cudaStreamDestroy(cudaStream_t s)//销毁流
//显性同步
cudaStreamSynchronize()//同步单个流：等待该流上的命令都完成
cudaDeviceSynchronize()//同步所有流同步：等待整个设备上流都完成
cudaStreamWaitEvent()//通过某个事件：等待某个事件结束后执行该流上的命令
cudaStreamQuery()//查询一个流任务是否完成
//回调
cudaStreamAddCallback()//在任何点插入回调函数
//优先级
cudaStreamCreateWithPriority()
cudaDeviceGetStreamPriorityRange()
3.例子
//创建两个流
cudaStream_t stream[2];
for (int i = 0; i < 2; ++i)
cudaStreamCreate(&stream[i]);
float* hostPtr;
cudaMallocHost(&hostPtr, 2 * size);
...
//两个流，每个流有三个命令
for (int i = 0; i < 2; ++i) {
//从主机内存复制数据到设备内存
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
//执行Kernel处理谁被内存
MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
//从设备内存到主机内存
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
...
//销毁流
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);
说明：
- 以上代码定义了两个流，每个流有三个命令，见注释
- 同一个流内顺序执行，流与流见异步执行
- 重叠行为参见[1]
4.其他问题
- 函数中若不指定流或者赋值为0，则为默认流
- 流还有隐性同步问题和重叠行为的问题
[1].http://docs.nvidia.com/cuda/cuda-c-programming-guide/#streams
