CUDA ---- Stream and Event

時間 2019-11-20

標籤 cuda stream event 欄目 Apache 简体版

原文原文鏈接

Streamhtml

通常來講，cuda c並行性表如今下面兩個層面上：編程

Kernel level
Grid level

到目前爲止，咱們討論的一直是kernel level的，也就是一個kernel或者一個task由許多thread並行的執行在GPU上。Stream的概念是相對於後者來講的，Grid level是指多個kernel在一個device上同時執行。api

Stream和event簡介

Cuda stream是指一堆異步的cuda操做，他們按照host代碼調用的順序執行在device上。Stream維護了這些操做的順序，並在全部預處理完成後容許這些操做進入工做隊列，同時也能夠對這些操做進行一些查詢操做。這些操做包括host到device的數據傳輸，launch kernel以及其餘的host發起由device執行的動做。這些操做的執行老是異步的，cuda runtime會決定這些操做合適的執行時機。咱們則可使用相應的cuda api來保證所取得結果是在全部操做完成後得到的。同一個stream裏的操做有嚴格的執行順序，不一樣的stream則沒有此限制。併發

因爲不一樣stream的操做是異步執行的，就能夠利用相互之間的協調來充分發揮資源的利用率。典型的cuda編程模式咱們已經熟知了：異步

將輸入數據從host轉移到device
在device上執行kernel
將結果從device上轉移回host

在許多狀況下，花費在執行kernel上的時間要比傳輸數據多得多，因此很容易想到將cpu和gpu之間的溝通時間隱藏在其餘kernel執行過程當中，咱們能夠將數據傳輸和kernel執行放在不一樣的stream中來實現此功能。Stream能夠用來實現pipeline和雙buffer（front-back）渲染。async

Cuda API可分爲同步和異步兩類，同步函數會阻塞host端的線程執行，異步函數會馬上將控制權返還給host從而繼續執行以後的動做。異步函數和stream是grid level並行的兩個基石。函數

從軟件角度來看，不一樣stream中的不一樣操做能夠並行執行，可是硬件角度卻不必定如此。這依賴於PCIe連接或者每一個SM可得到的資源，不一樣的stream仍然須要等待別的stream來完成執行。下面會簡單介紹在不一樣CC版本下，stream在device上的行爲。性能

Cuda Streams

全部的cuda操做（包括kernel執行和數據傳輸）都顯式或隱式的運行在stream中，stream也就兩種類型，分別是：測試

隱式聲明stream（NULL stream）
顯示聲明stream（non-NULL stream）

默認狀況下是NULL stream，在以前未涉及到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);

從device角度看，全部者三個操做都是使用的默認stream，而且按照代碼從上到下的順序依次執行，device自己是不知道其餘的host操做怎樣執行的。從host角度來看，數據傳輸都是同步的而且會一直等待，直到操做完成。不過不一樣於數據傳輸，Kernel的launch是異步的，host差很少馬上就能從新獲得控制權，不用管kernel是否執行完畢，從而進行下一步動做。很明顯，這種異步行爲有助於重疊device和host之間的運算時間。

上文內容在以前博文都有涉及，這裏特別說明的是數據傳輸，它也是能夠異步執行的，這就用到了本次講的stream，咱們必須顯示的聲明一個stream來分派它的執行。下面版本是異步版本的cudaMemcpy：

cudaError_t cudaMemcpyAsync(void* dst, const void* src, size_t count,cudaMemcpyKind kind, cudaStream_t stream = 0);

注意新增長的最後一個參數。這樣，在host issue了這個函數給device執行後，控制權能夠馬上返還給host。上面代碼使用了默認stream，若是要聲明一個新的stream則使用下面的API定義一個：

cudaError_t cudaStreamCreate(cudaStream_t* pStream);

這樣就定義了一個可使用在cuda異步API函數中stream。使用該函數的一個比較常見的錯誤，或者說容易引發混亂的地方是，這個函數返回的error code多是上一次調用異步函數產生的。也就是說，函數返回error並非調用該函數產生error的必要條件。

當執行一次異步數據傳輸時，咱們必須使用pinned（或者non-pageable）memory。Pinned memory的分配以下，具體請參見前面博文：

cudaError_t cudaMallocHost(void **ptr, size_t size);
cudaError_t cudaHostAlloc(void **pHost, size_t size, unsigned int flags);

經過在將該內存pin到host的虛擬內存上，就能夠將該memory的物理位置強制分配到CPU內存中以便使之在整個程序生命週期中保持不變。不然的話，操做系統可能會在任意時刻改變該host端的虛擬內存對應的物理地址。假設異步數據傳輸函數沒有使用pinned host memory的話，操做系統就可能將數據從一塊物理空間移動到另外一塊物理空間（由於是異步的，CPU在執行其餘的動做就可能影響這塊數據），而此時cuda runtime正在執行數據的傳輸，這會致使不肯定的行爲。

在執行kernel時要想設置stream的話，也是很簡單的，一樣只要加一個stream參數就好：

kernel_name<<<grid, block, sharedMemSize, stream>>>(argument list);

// 非默認的stream聲明
cudaStream_t stream;
// 初始化
cudaStreamCreate(&stream);
// 資源釋放
cudaError_t cudaStreamDestroy(cudaStream_t stream);

當執行資源釋放的時候，若是仍然有stream的工做沒幹完，那麼雖然該函數仍然會馬上返回，可是相關的工做作完後，這些資源纔會自動的釋放掉。

因爲全部stram的執行都是異步的，就須要一些API在必要的時候作同步操做：

cudaError_t cudaStreamSynchronize(cudaStream_t stream);
cudaError_t cudaStreamQuery(cudaStream_t stream);

第一個會強制host阻塞等待，直至stream中全部操做完成爲止；第二個會檢查stream中的操做是否所有完成，即便有操做沒完成也不會阻塞host。若是全部操做都完成了，則返回cudaSuccess，不然返回cudaErrorNotReady。

下面看一下一個代碼片斷來幫助理解：

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中。

上圖就跟流水線同樣差很少的道理，很少說。須要注意的是，上圖中數據傳輸的操做並非並行執行的，即便他們是在不一樣的stream中。按慣例，這種狀況確定就是硬件資源的鍋了，硬件資源就那麼些，軟件層面作的優化無非就是儘可能讓全部硬件資源一刻不停的被利用起來（萬惡的資本主義，嗯……），而這裏就是PCIe卡了瓶頸。固然從編程角度來看，這些操做依然是相互獨立的，只是他們要共享硬件資源，就不得不是串行的。有兩個PCIe就能夠重疊這兩次數據傳輸操做，不過也是要保證不一樣的stream和不一樣的傳輸方向。

最大併發kernel數目是依賴於device自己的，Fermi支持16路並行，Kepler是32。並行數是受限於shared memory，寄存器等device資源。

Stream Scheduling

概念上來講，全部stream是同時運行的。可是，事實上一般並不是如此。

False Dependencies

儘管Fermi最高支持16路並行，可是在物理上，全部stream是被塞進硬件上惟一一個工做隊列來調度的，當選中一個grid來執行時，runtime會查看task的依賴關係，若是當前task依賴前面的task，該task就會阻塞，因爲只有一個隊列，後面的都會跟着等待，即便後面的task是別的stream上的任務。就以下圖所示：

C和P以及R和X是能夠並行的，由於他們在不一樣的stream中，可是ABC，PQR以及XYZ卻不行，好比，在B沒完成以前，C和P都在等待。

Hyper-Q

僞依賴的狀況在Kepler系列裏獲得瞭解決，採用的一種叫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能夠搶佔低優先級執行。不過優先級屬性只對kernel有效，對數據傳輸無效。此外，若是設置的優先級超出了可設置範圍，則會自動設置成最高或者最低。有效可設置範圍可用下列函數查詢：

cudaError_t cudaDeviceGetStreamPriorityRange(int *leastPriority, int *greatestPriority);

顧名思義，leastPriority是下限，gretestPriority是上限。老規矩，數值較小則擁有較高優先級。若是device不支持優先級設置，則這兩個值都返回0。

Cuda Events

Event是stream相關的一個重要概念，其用來標記strean執行過程的某個特定的點。其主要用途是：

同步stream執行
操控device運行步調

Cuda api提供了相關函數來插入event到stream中和查詢該event是否完成（或者叫知足條件？）。只有當該event標記的stream位置的全部操做都被執行完畢，該event纔算完成。關聯到默認stream上的event則對全部的stream有效。

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

Events標記了stream執行過程當中的一個點，咱們就能夠檢查正在執行的stream中的操做是否到達該點，咱們能夠把event當成一個操做插入到stream中的衆多操做中，當執行到該操做時，所作工做就是設置CPU的一個flag來標記表示完成。下面函數將event關聯到指定stream。

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上，可是要注意，若是兩者任意一個關聯到了non-NULL stream上，時間間隔可能要比指望的大。這是由於cudaEventRecord是異步發生的，咱們沒辦法保證度量出來的時間剛好就是兩個event之間，因此只是想要gpu工做的時間間隔，則stop和strat都關聯到默認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能夠被分爲同步（NULL stream）和異步（non-NULL stream）兩種，同步異步是針對host來說的，異步stream不會阻塞host的執行，而大多數同步stream則會阻塞host，不過kernel launch例外，不會阻塞host。

此外，異步stream又能夠被分爲阻塞和非阻塞兩種，阻塞非阻塞是異步stream針對同步stream來說的。異步stream若是是阻塞stream，那麼同步stream會阻塞該異步stream中的操做。若是異步stream是非阻塞stream，那麼該stream不會阻塞同步stream中的操做（有點繞……）。

阻塞和非阻塞stream

使用cudaStreamCreate建立的是阻塞stream，也就是說，該stream中執行的操做會被早先執行的同步stream阻塞。一般來講，當issue一個NULL stream時，cuda context會等待以前全部阻塞stream完成後才執行該NULL stream，固然全部阻塞stream也會等待以前的NULL stream完成纔開始執行。

例如：

kernel_1<<<1, 1, 0, stream_1>>>();
kernel_2<<<1, 1>>>();
kernel_3<<<1, 1, 0, stream_2>>>();

從device角度來講，這三個kernel是串行依次執行的，固然從host角度來講，倒是並行非阻塞的。除了經過cudaStreamCreate生成的阻塞stream外，咱們還能夠經過下面的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)

若是以前的kernel_1和kernel_3的stream被定義成第二種，就不會被阻塞。

Implicit Synchronization

Cuda有兩種類型的host和device之間同步：顯式和隱式。咱們以前已經瞭解到顯式同步API有：

cudaDeviceSynchronize
cudaStreamSynchronize
cudaEventSynchronize

這三個函數由host顯式的調用，在device上執行。

隱式同步咱們也瞭解過，好比cudaMemcpy就會隱式的同步device和host，由於該函數同步做用只是數據傳輸的反作用，因此稱爲隱式。瞭解這些隱式同步是很中要的，由於不經意的調用這樣一個函數可能會致使性能急劇下降。

隱式同步是cuda編程中比較特殊狀況，由於隱式同步行爲可能會致使意外的阻塞行爲，一般發生在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上的全部操做。該函數會致使host等待全部device上的運算或者數據傳輸操做完成。顯而易見，該函數是個heavyweight的函數，咱們應該儘可能減小這類函數的使用。

經過使用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

cudaEventBlockingSync說明該event會阻塞host。cudaEventSynchronize默認行爲是使用CPU時鐘來固定的查詢event狀態。使用cudaEventBlockingSync，調用線程會進入休眠，將控制權交給其餘線程或者進程，直到event完成爲止。可是這樣會致使少許的CPU時鐘浪費，也會增長event完成和喚醒線程的之間的時間消耗。

cudaEventDisableTiming指定event只能用來同步，而且不須要記錄計時數據。這樣扔掉記錄時間戳的消耗能夠提升cuudaStreamWaitEvent和cudaEventQuery的調用性能。

cudaEventInterprocess指定event能夠被用來做爲inter-process event。

NVIDIA CUDA板塊：https://developer.nvidia.com/cuda-zone

CUDA在線文檔：http://docs.nvidia.com/cuda/index.html#

轉載原文註明：http://www.cnblogs.com/1024incn/p/5891051.html

填坑中~~