CUDA之Dynamic Parallelism詳解（二）

Bruce_0712發表於2017-03-19
Parallel

CUDA 5.0中引入動態並行化，使得在device端執行的kernel的執行緒也能跟在host上一樣launch kernels，只有支援CC3.5或者以上的裝置中才能支援。動態並行化使用CUDA Device Runtime library（cudadevrt），它是一個能在device code中呼叫的CUDA runtime子集。

編譯連結

為了支援動態並行化，必須使用兩步分離編譯和連結的過程：首先，設定-c和-rdc=true（–relocatable-device-code=true）來生成relocatable device code來進行後續連結，可以使用-dc(–device -c)來合併這兩個選項；然後將上一步目標檔案和cudadevrt庫進行連線生成可執行檔案，-lcudadevrt。過程如下圖

或者簡化成一步

執行、同步

在CUDA程式設計模型中，一組執行的kernel的執行緒塊叫做一個grid。在CUDA動態並行化，parent grid能夠呼叫child grids。child grid繼承parant grid的特定屬性和限制，如L1 cache、shared_memory、棧大小。如果一個parent grid有M個block和N個thread，如果對child kernel launch沒有控制的話，那個將產生M*N個child kernel launch。如果想一個block產生一個child kernel，那麼只需要其中一個執行緒launch a kernel就行。如下

grid lanuch是完全巢狀的，child grids總是在發起它們的parent grids結束前完成，這可以看作是一個一種隱式的同步。

如果parent kernel需要使用child kernel的計算結果，也可以使用CudaDeviceSynchronize(void)進行顯示的同步，這個函式會等待一個執行緒塊發起的所有子kernel結束。往往不知道一個執行緒塊中哪些子kernel已經執行，可以通過下述方式進行一個執行緒塊級別的同步

CudaDeviceSynchronize(void)呼叫開銷較大，不是必須的時候，儘量減少使用，同時不要在父kernel退出時呼叫，因為結束時存在上述介紹的隱式同步。

記憶體一致

當子 grids開始與結束之間，父grids和子grids有完全一致的global memory view。

當子kernel launch的時候，global memory檢視不一致。

在子kernel launch之後，顯示同步之前，parent grid不能對 child grid讀取的記憶體做寫入操作，否則會造成race condition。

向Child grids傳遞指標

指標的傳遞存在限制：

  • 可以傳遞的指標：global memory（包括__device__變數和malloc分配的記憶體），zero-copy host端記憶體，常量記憶體。
  • 不可以傳遞的指標：shared_memory（__shared__變數）, local memory（包括stack變數）

Device Streams和Events

所有在device上建立的streams都是non-blocking的，不支援預設NULL stream的隱式同步。建立流的方式如下

一旦一個device stream被建立，它能被一個執行緒塊中其他執行緒使用。只有當這個執行緒塊完成執行的時候，這個stream才能被其他執行緒塊或者host使用。反之亦然。

Event也是支援的，不過有限制，只支援在不同stream之間使用cudaStreamWaitEvent()指定執行順序，而不能使用event來計時或者同步。

Recursion Depth和Device Limits

遞迴深度包括兩個概念：

  • nesting depth：遞迴grids的最大巢狀層次，host端的為0；
  • synchronization depth：cudaDeviceSynchronize()能呼叫的最大巢狀層次，host端為1，cudaLimitDevRuntimeSyncDepth應該設定為maximum 所以你吃肉你咋體on depth加1，設定方式如 cudaDeviceLimit(cudaLimitDevRuntimeSyncDepth, 4).

maximum nesting depth有硬體限制，在CC3.5中， 對depth 的限制為24. synchronization depth也一樣。

從外到內，直到最大同步深度，每一次層會保留一部分記憶體來儲存父block的上下文資料，即使這些記憶體沒有被使用。所以遞迴深度的設定需要考慮到每一層所預留的記憶體。

另外還有一個限制是待處理的子grid數量。pending launch buffer用來維持launch queue和追蹤當前執行kernel的狀態。通過

來設定合適的限制。否則通過cudaGetLastError()呼叫可以返回CudaErrorLaunchPendingCountExceeded的錯誤。

動態並行化執行有點類似樹的結構，但與CPU上樹處理也有些不同。類似深度小，分支多，比較茂密的樹的執行結構，比較適合動態並行化的處理。深度大，每層節點少的樹的執行結構，則不適合動態並行化。

characteristic tree processing dynamic parallelism
node thin (1 thread) thick (many threads)
branch degree small (usually < 10) large (usually > 100)
depth large small


相關文章