CUDA 網格級併發流 2

上節中描述到cuda中可以通過流來實現網格級併發，按照流的劃分主要分為兩種：

空流為乙個同步流，在空流中的大部分操作都會堵塞主機，kernel執行除外，

而非空流為乙個非同步流，其上所有的操作都不堵塞主機執行。

而非空流按照與空流直接的關係還可分為兩種型別

非空流中的堵塞和非堵塞主要是針對與空流之前的關係而劃分，下面來根據幾種情況來說明：

主機上非空流是非同步流，其上所有的操作都不會阻塞主機執行。相應地，隱式的空流是同步流，大多數新增到空流上的操作都會導致主機在先前所有的操作產生阻塞。

下面以乙個例子來說明上述情況：

const int n =3;
cudastream_t *stream = (cudastream_t*)malloc(sizeof(cudastream_t)*n);
for (int i = 0; i < n; i++)
kernel1 << <1, 1>> > (); //空流 會堵塞kernel 2和kernel 3
kernel2 << <1,1,0,stream[0]>> > (); //非空流+阻塞流
kernel3 << <1, 1, 0, stream[1] >> > ();//非空流+阻塞流
for (int i = 0; i < n; i++)

上述例子中kernel 1為空流中的操作， kernel2為stream[0], kernel3為stream[1], 上述三個操作對主機而言都是非同步的，所有需要進行同步操作，當kernel分別載入到流中，在gpu真正的執行順序並不是我們想象的三個kernel同時並行，而是先執行kernel1, kernel1執行完後才會同時執行kernel 2和3，這是因為空流對於非空流的操作具有乙個堵塞作用。非空流的操作可以被空流中的操作所堵塞，即kerne1執行完後，才會執行kernel 2和3.操作被發布到阻塞流中時，在操作執行之前，都會被掛起，直至空流中的操作完成。上句話的意思是當kernel2執行之前，會一直等待kernel 1執行完畢之後才會執行。

非空堵塞流中的操作會堵塞空流中的操作。

將上述用例進行稍微改造：

const int n =3;
cudastream_t *stream = (cudastream_t*)malloc(sizeof(cudastream_t)*n);
for (int i = 0; i < n; i++)
kernel1 << <1,1,0,stream[0]>> > (); //非空流+阻塞流，
kernel2 << <1, 1>> > (); //空流 會堵塞kernel 3
kernel3 << <1, 1, 0, stream[1] >> > ();//非空流+阻塞流
for (int i = 0; i < n; i++)

其中kernel1操作加入到非空流stream[0]中， kernel為乙個空流操作，kernel3為乙個非空流stream[1]中，由於非空堵塞流中的操作會堵塞空流中的操作，所以kernel2會一直等待kernel1執行完畢之後，才會得到執行，而kernel2為空流，同時會堵塞kernel3,所以執行順序為kernel1->kernel 2->kernel 3序列執行的效果。

當操作被發布到空流中時，在操作被執行之前，cuda上下文會等待所有先前操作發布到所有的阻塞流中。

這句話的意思是kernel 2在執行之前，會一直等待之前的非空堵塞流操作執行完畢，即kernel1執行完畢才會執行kernel2.

非空非堵塞流不會對空流操作堵塞操作，空流也不會堵塞非空非堵塞流。

cuda中使用cudastreamcreate ()api建立的是非空堵塞流，如果要建立非空非堵塞流需要使用到如下api:

__host__ __device__ cudaerror_t cudastreamcreatewithflags ( cudastream_t* pstream, unsigned int flags )

其中flags為建立的流型別：

cudastreamdefault：預設值，建立流為非空堵塞流，會對空流造成堵塞。

cudastreamnonblocking：建立非空非堵塞流，不會對空流造成堵塞。

下面非乙個非空非堵塞流的例子：

const int n = 3;
cudastream_t *stream = (cudastream_t*)malloc(sizeof(cudastream_t)*n);
for (int i = 0; i < n; i++)
kernel1 << <1,1,0,stream[0]>> > ();
kernel2 << <1, 1>> > ();
kernel3 << <1, 1, 0, stream[1] >> > ();
for (int i = 0; i < n; i++)

上述例子kernel 1和kernel 3為乙個非空流， kernel 1和kernel 2,kernel 3為並行執行。

非空流之間不管是堵塞還是非堵塞，都不會相互影響，在上章節中的例子：

or (int i= 0; i< nstreams;i++)
for(int i =0; i< nstream;i++)

即使建立的流為非空流，相互之前也不會受影響。

CUDA 網格級併發流 2

CUDA 學習（六）執行緒網格

WinForm畫網格並填充顏色

CUDA中二級指標表示二級陣列

CUDA 網格級併發 流 2

CUDA 學習（六） 執行緒網格

WinForm畫網格並填充顏色

CUDA中二級指標表示二級陣列

相關推薦

CUDA 網格級併發流 2

CUDA 學習（六）執行緒網格