天天看點

CUDA實踐指南(二十三)

線程和塊啟發式:

中等優先級:每塊的線程數應該是32個線程的倍數,因為這提供了最優的計算效率并促進了合并。

每個網格的塊的尺寸和大小以及每個塊的線程的尺寸和大小都是重要的因素。 這些參數的多元方面允許更輕松地将多元問題映射到CUDA,并且不影響性能。 是以,本節讨論尺寸而不是次元。

選擇執行配置參數應該一起完成; 然而,有一些啟發式方法适用于每個參數。 在選擇第一個執行配置參數時(每個網格的塊數或網格大小),主要關心的是讓整個GPU保持繁忙。 網格中的塊數應該大于多處理器的數量,以便所有多處理器都具有

至少要執行一個塊。 此外,每個多處理器應該有多個活動塊,以便不等待<code>__syncthreads()</code>的塊可以保持硬體繁忙。 這項建議取決于資源的可用性; 是以,應該在第二個執行參數 - 每個塊的線程數或塊大小 - 以及共享記憶體使用情況下确定。 為了擴充到未來的裝置,每個核心啟動的塊數應該是數千。

選擇塊大小時,務必記住多個并發塊可以駐留在多處理器上,是以占用率不僅僅取決于塊大小。 特别是,更大的塊大小并不意味着更高的占用率。 例如,在計算能力為1.1或更低的裝置上,最大塊大小為512線程的核心導緻66%的占用率,因為此類裝置上每個多處理器的最大線程數為768.是以,隻有一個 塊可以為每個多處理器激活。 但是,在這種裝置上每個塊有256個線程的核心可以使三個駐留的活動塊占用100%。

如占有率所述,較高的占有率并不總是等同于更好的表現。 例如,将占有率從66%提高到100%通常不會轉化為性能類似的提高。 較低的占用核心每個線程會擁有比較高的占用核心更多的寄存器,這可能導緻寄存器溢出到本地記憶體的數量減少。 通常情況下,一旦達到50%的占有率,占有率的增加并不意味着改善的性能。 在某些情況下,可以通過更少的翹曲完全覆寫延遲,特别是通過指令級并行(ILP)。

選擇塊大小涉及很多這樣的因素,并且不可避免地需要一些實驗。 但是,應遵循一些經驗法則:

每個塊的線程數應該是warp大小的倍數,以避免在填充不足的warp上浪費計算并促進合并。

每塊應該使用至少64個線程,并且每個多處理器隻有多個并發塊。

每塊128到256個線程是更好的選擇,并且對于不同塊大小的實驗來說是一個很好的初始範圍。

如果延遲影響性能,則每個多處理器使用幾個(3到4個)較小的線程塊,而不是一個較大的線程塊。 這對經常調用__syncthreads()的核心特别有利。

請注意,當一個線程塊配置設定的寄存器比多處理器上可用的寄存器多時,核心啟動失敗,因為當請求太多的共享記憶體或太多線程時,核心啟動失敗。

繼續閱讀