本文從CSDN上轉(zhuǎn)移過來：
http://blog.csdn.net/mounty_fsc/article/details/51092933
1.流（stream）的理解
一個流對應(yīng)并發(fā)的概念，是一組順序執(zhí)行的操作（可能由多個主機線程發(fā)出）；
多個流對應(yīng)并行的概念，因為發(fā)生順序具有不確定性。
2.相關(guān)函數(shù)
//基本函數(shù)
cudaStream_t stream//定義流
cudaStreamCreate(cudaStream_t * s)//創(chuàng)建流
cudaStreamDestroy(cudaStream_t s)//銷毀流
//顯性同步
cudaStreamSynchronize()//同步單個流：等待該流上的命令都完成
cudaDeviceSynchronize()//同步所有流同步：等待整個設(shè)備上流都完成
cudaStreamWaitEvent()//通過某個事件：等待某個事件結(jié)束后執(zhí)行該流上的命令
cudaStreamQuery()//查詢一個流任務(wù)是否完成
//回調(diào)
cudaStreamAddCallback()//在任何點插入回調(diào)函數(shù)
//優(yōu)先級
cudaStreamCreateWithPriority()
cudaDeviceGetStreamPriorityRange()
3.例子
//創(chuàng)建兩個流
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) {
//從主機內(nèi)存復(fù)制數(shù)據(jù)到設(shè)備內(nèi)存
cudaMemcpyAsync(inputDevPtr + i * size, hostPtr + i * size, size, cudaMemcpyHostToDevice, stream[i]);
//執(zhí)行Kernel處理誰被內(nèi)存
MyKernel <<<100, 512, 0, stream[i]>>>(outputDevPtr + i * size, inputDevPtr + i * size, size);
//從設(shè)備內(nèi)存到主機內(nèi)存
cudaMemcpyAsync(hostPtr + i * size, outputDevPtr + i * size, size, cudaMemcpyDeviceToHost, stream[i]);
}
...
//銷毀流
for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);
說明：
- 以上代碼定義了兩個流，每個流有三個命令，見注釋
- 同一個流內(nèi)順序執(zhí)行，流與流見異步執(zhí)行
- 重疊行為參見[1]
4.其他問題
- 函數(shù)中若不指定流或者賦值為0，則為默認(rèn)流
- 流還有隱性同步問題和重疊行為的問題
[1].http://docs.nvidia.com/cuda/cuda-c-programming-guide/#streams
