如何設定CUDA Kernel中的grid_size和block_size？

語言: CN / TW / HK
時間 2021-12-09 17:01:54
華為雲Stack Iass首席架構師聊設計架構>>>

 

撰文 | 柳俊丞

 

導語：在剛接觸 CUDA 程式設計時，很多人都會疑惑在啟動一個kernel時，三個尖括號裡面的引數應該如何設定？這些引數受到哪些因素的制約？以及他們如何影響 kernel 執行的效能？本文參考 CUDA 官方文件，分析了這些引數應該如何設定。

 

一般而言，我們在程式碼中會看到使用以下方式啟動一個 CUDA kernel：

 

cuda_kernel<<<grid_size, block_size,  0, stream >>>(...)

 

cuda_kernel  是 global function 的標識， (...)  中是呼叫 cuda_kernel 對應的引數，這兩者和 C++ 的語法是一樣的，而  <<<grid_size, block_size, 0, stream>>>  是 CUDA 對 C++ 的擴充套件，稱之為 Execution Configuration（ http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#execution-configuration ），參考 CUDA C++ Programming Guide ( http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#abstract ，後續簡稱 Guide ) 中的介紹：

The execution configuration is specified by inserting an expression of the form <<< Dg, Db, Ns, S >>> between the function name and the parenthesized argument list, where:

  • Dg is of type dim3 (see  dim3 ) and specifies the dimension and size of the grid, such that Dg.x * Dg.y * Dg.z equals the number of blocks being launched;

  • Db is of type dim3 (see  dim3 ) and specifies the dimension and size of each block, such that Db.x * Db.y * Db.z equals the number of threads per block;

  • Ns is of type size_t and specifies the number of bytes in shared memory that is dynamically allocated per block for this call in addition to the statically allocated memory; this dynamically allocated memory is used by any of the variables declared as an external array as mentioned in  shared ; Ns is an optional argument which defaults to 0;

  • S is of type cudaStream_t and specifies the associated stream; S is an optional argument which defaults to 0.

 

Dg 代表的是 grid 的維度，Db 代表 block 的維度，型別為 dim3，如果是簡單的一維結構，也就是除了x以外，yz兩個維度對應的值都是1，Dg 和 Db 也可以直接用 x 維度對應的數字代替，也就是文章一開始的表示方式，對 grid dim 與 block dim 兩者更具體的說明可以參考 Programming Model（ http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#programming-model ）

 

接下來我們討論一下這兩個值一般應該取什麼值。

grid_size 和 block_size 分別代表了本次 kernel 啟動對應的 block 數量和每個 block 中 thread 的數量，所以顯然兩者都要大於 0。

 

Guide 中 K.1. Features and Technical Specifications （ http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#features-and-technical-specifications ）指出， Maximum number of threads per block 以及 Maximum x- or y-dimension of a block 都是 1024，所以 block_size 最大可以取 1024。

 

同一個 block 中，連續的 32 個執行緒組成一個 warp，這 32 個執行緒每次執行同一條指令，也就是所謂的 SIMT，即使最後一個 warp 中有效的執行緒數量不足 32，也要使用相同的硬體資源，所以 block_size 最好是 32 的整數倍。

block 有時也會被稱之為  Cooperative Thread Arrays （ http://docs.nvidia.com/cuda/parallel-thread-execution/index.html#cooperative-thread-arrays ），參考

The Parallel Thread Execution (PTX) programming model is explicitly parallel: a PTX program specifies the execution of a given thread of a parallel thread array. A cooperative thread array, or CTA, is an array of threads that execute a kernel concurrently or in parallel.

 

Threads within a CTA can communicate with each other. To coordinate the communication of the threads within the CTA, one can specify synchronization points where threads wait until all threads in the CTA have arrived.

與 block 對應的硬體級別為 SM，SM 為同一個 block 中的執行緒提供通訊和同步等所需的硬體資源，跨 SM 不支援對應的通訊，所以一個 block 中的所有執行緒都是執行在同一個 SM 上的，而且因為執行緒之間可能同步，所以一旦 block 開始在 SM 上執行，block 中的所有執行緒同時在同一個 SM 中執行（併發，不是並行），也就是說 block 排程到 SM 的過程是原子的。SM 允許多於一個 block 在其上併發執行，如果一個 SM 空閒的資源滿足一個 block 的執行，那麼這個 block 就可以被立即排程到該 SM 上執行，具體的硬體資源一般包括暫存器、shared memory、以及各種排程相關的資源。

 

這裡的排程相關的資源一般會表現為兩個具體的限制， Maximum number of resident blocks per SM 和 Maximum number of resident threads per SM ，也就是 SM 上最大同時執行的 block 數量和執行緒數量。因為 GPU 的特點是高吞吐高延遲，就像一個自動扶梯一分鐘可以運送六十個人到另一層樓，但是一個人一秒鐘無法通過自動扶梯到另一層樓，要達到自動扶梯可以運送足夠多的人的目標，就要保證扶梯上同一時間有足夠多的人，對應到 GPU，就是要儘量保證同一時間流水線上有足夠多的指令。

 

要到達這個目的有多種方法，其中一個最簡單的方法是讓儘量多的執行緒同時在 SM 上執行，SM 上併發執行的執行緒數和SM 上最大支援的執行緒數的比值，被稱為 Occupancy，更高的 Occupancy 代表潛在更高的效能。

 

顯然，一個 kernel 的 block_size 應大於 SM 上最大執行緒數和最大 block 數量的比值，否則就無法達到 100% 的 Occupancy，對應不同的架構，這個比值不相同，對於 V100 、 A100、 GTX 1080 Ti 是 2048 / 32 = 64，對於 RTX 3090 是 1536 / 16 = 96，所以為了適配主流架構，如果靜態設定 block_size 不應小於 96。考慮到 block 排程的原子性，那麼 block_size 應為 SM 最大執行緒數的約數，否則也無法達到 100% 的 Occupancy，主流架構的 GPU 的 SM 最大執行緒數的公約是 512，96 以上的約數還包括 128 和 256，也就是到目前為止，block_size 的可選值僅剩下 128 / 256 / 512 三個值。

 

還是因為 block 排程到 SM 是原子性的，所以 SM 必須滿足至少一個 block 執行所需的資源，資源包括 shared memory 和暫存器， shared memory 一般都是開發者顯式控制的，而如果 block 中執行緒的數量 * 每個執行緒所需的暫存器數量大於 SM 支援的每 block 暫存器最大數量，kernel 就會啟動失敗。

 

目前主流架構上，SM 支援的每 block 暫存器最大數量為 32K 或 64K 個 32bit 暫存器，每個執行緒最大可使用 255 個 32bit 暫存器，編譯器也不會為執行緒分配更多的暫存器，所以從暫存器的角度來說，每個 SM 至少可以支援 128 或者 256 個執行緒，block_size 為 128 可以杜絕因暫存器數量導致的啟動失敗，但是很少的 kernel 可以用到這麼多的暫存器，同時 SM 上只同時執行 128 或者 256 個執行緒，也可能會有潛在的效能問題。但把 block_size 設定為 128，相對於 256 和 512 也沒有什麼損失，128 作為 block_size 的一個通用值是非常合適的。

 

確定了 block_size 之後便可以進一步確定 grid_size，也就是確定總的執行緒數量，對於一般的 elementwise kernel 來說，總的執行緒數量應不大於總的 element 數量，也就是一個執行緒至少處理一個 element，同時 grid_size 也有上限，為 Maximum x-dimension of a grid of thread blocks ，目前在主流架構上都是 2^31 - 1，對於很多情況都是足夠大的值。

 

有時為每個 element 建立一個執行緒是可行的，因為執行緒的建立在 GPU 上是一個開銷足夠低的操作，但是如果每個執行緒都包含一個公共的操作，那麼執行緒數的增多，也代表著這部分的開銷變大，比如

__ global__
void kernel(const float* x, const float* v, float* y) {
    const  float sqrt_v =  sqrt(*v);
    const  int idx = blockIdx* gridDim.x + threadIdx.x;
   y[idx] = x[idx] * sqrt_v;
}

 

這個 kernel 中對 v 的處理是公共的，如果我們減少執行緒的數量並迴圈處理 y 和 x，那麼 sqrt(*v)  的開銷就會相應降低，但是 grid_size 的數值不應低於 GPU 上 SM 的數量，否則會有 SM 處於空閒狀態。

 

我們可以想象，GPU 一次可以排程 SM 數量 * 每個 SM 最大 block 數個 block，因為每個 block 的計算量相等，所以所有 SM 應幾乎同時完成這些 block 的計算，然後處理下一批，這其中的每一批被稱之為一個 wave。想象如果 grid_size 恰好比一個 wave 多出一個 block，因為 stream 上的下個 kernel 要等這個 kernel 完全執行完成後才能開始執行，所以第一個 wave 完成後，GPU 上將只有一個 block 在執行，GPU 的實際利用率會很低，這種情況被稱之為 tail effect。

 

我們應儘量避免這種情況，將 grid_size 設定為精確的一個 wave 可能也無法避免 tail effect，因為 GPU 可能不是被當前 stream 獨佔的，常見的如 NCCL 執行時會佔用一些 SM。所以無特殊情況，可以將 grid_size 設定為數量足夠多的整數個 wave，往往會取得比較理想的結果，如果數量足夠多，不是整數個 wave 往往影響也不大。

 

綜上所述，普通的 elementwise kernel 或者近似的情形中，block_size 設定為 128，grid_size 設定為可以滿足足夠多的 wave 就可以得到一個比較好的結果了。但更復雜的情況還要具體問題具體分析，比如，如果因為 shared_memory 的限制導致一個 SM 只能同時執行很少的 block，那麼增加 block_size 有機會提高效能，如果 kernel 中有執行緒間同步，那麼過大的 block_size 會導致實際的 SM 利用率降低，這些我們有機會單獨討論。

 

其他人都在看

歡迎下載體驗OneFlow新一代開源深度學習框架：http://github.com/Oneflow-Inc/oneflow/

 

本文分享自微信公眾號 - OneFlow（OneFlowTechnology）。
如有侵權，請聯絡 [email protected] 刪除。
本文參與“OSC源創計劃”，歡迎正在閱讀的你也加入，一起分享。

「」
「」