Stream
一般來說，cuda c並行性表現在下面兩個層面上：
- Kernel level
- Grid level
Stream和event簡介
Cuda stream是指一堆異步的cuda操作，他們按照host代碼調用的順序執行在device上。
典型的cuda編程模式我們已經熟知了：
- 將輸入數據從host轉移到device
- 在device上執行kernel
- 將結果從device上轉移回host
Cuda Streams
所有的cuda操作（包括kernel執行和數據傳輸）都顯式或隱式的運行在stream中，stream也就兩種類型，分別是：
- 隱式聲明stream（NULL stream）
- 顯示聲明stream（non-NULL stream）
異步且基於stream的kernel執行和數據傳輸能夠實現以下幾種類型的並行：
- Host運算操作和device運算操作並行
- Host運算操作和host到device的數據傳輸並行
- Host到device的數據傳輸和device運算操作並行
- Device內的運算並行
下面代碼是常見的使用形式，默認使用NULL stream:
cudaMemcpy(..., cudaMemcpyHostToDevice);
kernel<<<grid, block>>>(...);
cudaMemcpy(..., cudaMemcpyDeviceToHost);
下面版本是異步版本的cudaMemcpy：
cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);
上面代碼使用了默認stream，如果要聲明一個新的stream則使用下面的API定義一個：
cudaError_t cudaStreamCreate(cudaStream_t* pStream);
Pinned memory的分配如下：
cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);
在執行kernel時要想設置stream的話，只要加一個stream參數就好：
kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);
// 非默認的stream聲明
cudaStream_t stream;
// 初始化
cudaStreamCreate(&stream);
// 資源釋放
cudaError_t cudaStreamDestroy(cudaStream_t stream);
所有stram的執行都是異步的，需要一些API在必要的時候做同步操作：
cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);
看一下代碼片段：
for (int i = 0; i < nStreams; i++) {
int offset = i * bytesPerStream;
cudaMemcpyAsync(&d_a[offset], &a[offset], bytePerStream, streams[i]);
kernel<<grid, block, 0, streams[i]>>(&d_a[offset]);
cudaMemcpyAsync(&a[offset], &d_a[offset], bytesPerStream, streams[i]);
}
for (int i = 0; i < nStreams; i++) {
cudaStreamSynchronize(streams[i]);
}
使用了三個stream，數據傳輸和kernel運算都被分配在了這幾個並發的stream中。
kernel數目是依賴於device本身的，Fermi支持16路並行，Kepler是32。並行數是受限於shared memory，寄存器等device資源。
Stream Scheduling
C和P以及R和X是可以並行的，因為他們在不同的stream中，但是ABC，PQR以及XYZ卻不行，比如，在B沒完成之前，C和P都在等待。
Hyper-Q
Hyper-Q的技術， Kepler上出現了32個工作隊列。實現了TPC上可以同時運行compute和graphic的應用。當然，如果超過32個stream被創建了，依然會出現偽依賴的情況。
Stream Priorities
對於CC3.5及以上版本，stream可以有優先級的屬性：
cudaError_t cudaStreamCreateWithPriority(cudaStream_t* pStream, unsigned int flags, int priority);
該函數創建一個stream，賦予priority的優先級，高優先級的grid可以搶占低優先級執行。
cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);
leastPriority是下限，gretestPriority是上限。數值較小則擁有較高優先級。如
Cuda Events
Event是stream用來標記strean執行過程的某個特定的點。其主要用途是：
- 同步stream執行
- 操控device運行步調
Creation and Destruction
// 聲明
cudaEvent_t event;
// 創建
cudaError_t cudaEventCreate(cudaEvent_t* event);
// 銷毀
cudaError_t cudaEventDestroy(cudaEvent_t event);
streeam的釋放，在操作完成后自動釋放資源。
Recording Events and Mesuring Elapsed Time
cudaError_t cudaEventRecord(cudaEvent_t event, cudaStream_t stream = 0);
等待event會阻塞調用host線程，同步操作調用下面的函數：
cudaError_t cudaEventSynchronize(cudaEvent_t event);
類似於cudaStreamSynchronize，等待event而不是整個stream執行完畢。使用API來測試event是否完成，該函數不會阻塞host：
cudaError_t cudaEventQuery(cudaEvent_t event);
該函數類似cudaStreamQuery。此外，還有專門的API可以度量兩個event之間的時間間隔：
cudaError_t cudaEventElapsedTime(float* ms, cudaEvent_t start, cudaEvent_t stop);
返回start和stop之間的時間間隔，單位是毫秒。Start和stop不必關聯到同一個stream上。
下面代碼簡單展示了如何使用event來度量時間：
// create two events
cudaEvent_t start, stop;
cudaEventCreate(&start);
cudaEventCreate(&stop);
// record start event on the default stream
cudaEventRecord(start);
// execute kernel
kernel<<<grid, block>>>(arguments);
// record stop event on the default stream
cudaEventRecord(stop);
// wait until the stop event completes
cudaEventSynchronize(stop);
// calculate the elapsed time between two events
float time;
cudaEventElapsedTime(&time, start, stop);
// clean up the two events
cudaEventDestroy(start);
cudaEventDestroy(stop);
Stream Synchronization
由於所有non-default stream的操作對於host來說都是非阻塞的，就需要相應的同步操作。
從host的角度來看，cuda操作可以被分為兩類：
- Memory相關的操作
- Kernel launch
Kernel launch對於host來說都是異步的，許多memory操作則是同步的，比如cudaMemcpy，cuda runtime也會提供異步函數來執行memory操作。
阻塞和非阻塞stream
使用cudaStreamCreate創建的是阻塞stream，也就是說，該stream中執行的操作會被早先執行的同步stream阻塞。
例如：
kernel_1<<<1, 1, 0, stream_1>>>();
kernel_2<<<1, 1>>>();
kernel_3<<<1, 1, 0, stream_2>>>();
可以通過下面的API配置生成非阻塞stream：
cudaError_t cudaStreamCreateWithFlags(cudaStream_t* pStream, unsigned int flags);
// flag為以下兩種，默認為第一種，非阻塞便是第二種。
cudaStreamDefault: default stream creation flag (blocking)
cudaStreamNonBlocking: asynchronous stream creation flag (non-blocking)
Implicit Synchronization
Cuda有兩種類型的host和device之間同步：顯式和隱式。已經了解到顯式同步API有：
- cudaDeviceSynchronize
- cudaStreamSynchronize
- cudaEventSynchronize
這三個函數由host顯式的調用，在device上執行。
許多memory相關的操作都會影響當前device的操作，比如：
- A page-locked host memory allocation
- A device memory allocation
- A device memset
- A memory copy between two addresses on the same device
- A modification to the L1/shared memory confi guration
Explicit Synchronization
從grid level來看顯式同步方式，有如下幾種：
- Synchronizing the device
- Synchronizing a stream
- Synchronizing an event in a stream
- Synchronizing across streams using an event
可以使用cudaDeviceSynchronize來同步該device上的所有操作。通過使用cudaStreamSynchronize可以使host等待特定stream中的操作全部完成或者使用非阻塞版本的cudaStreamQuery來測試是否完成。
Cuda event可以用來實現更細粒度的阻塞和同步，相關函數為cudaEventSynchronize和cudaEventSynchronize，用法類似stream相關的函數。此外，cudaStreamWaitEvent提供了一種靈活的方式來引入stream之間的依賴關系：
cudaError_t cudaStreamWaitEvent(cudaStream_t stream, cudaEvent_t event);
該函數會指定該stream等待特定的event，該event可以關聯到相同或者不同的stream，對於不同stream的情況，如下圖所示：
Stream2會等待stream1中的event完成后繼續執行。
Configurable Events
Event的配置可用下面函數：
cudaError_t cudaEventCreateWithFlags(cudaEvent_t* event, unsigned int flags);
cudaEventDefault
cudaEventBlockingSync
cudaEventDisableTiming
cudaEventInterprocess