CUDA分支優化

標簽： cuda分支優化warpSM 2015-07-16 10:24 293人閱讀 評論(0) 收藏 舉報 分類： CUDA（26）

版權聲明：本文為博主原創文章，未經博主允許不得轉載。

在CUDA中，分支會極大的減弱性能，因為ＳＭ沒有分支預測，因此只能讓束內線程在每個分支上都執行一遍，當然如果某個分支沒有線程執行，就可以忽略，因此要減少分支的數目?？梢院唵蔚恼f：
1. 同一個warp中的所有線程執行相同的命令。
2. 并不是所有線程都會執行。

產生分支的一個常見場景就是if,else語句的使用，比如常用的邊界判定。下面從一個實力說起。看如下代碼：

<code class="hljs cs has-numbering"> <span class="hljs-keyword">int</span> tid = threadIdx.x;<span class="hljs-keyword">if</span>(tid == <span class="hljs-number">0</span>)<span class="hljs-keyword">var</span> = <span class="hljs-keyword">var</span> + <span class="hljs-number">1</span>;<span class="hljs-keyword">else</span><span class="hljs-keyword">var</span> = <span class="hljs-keyword">var</span> + <span class="hljs-number">2</span>;<span class="hljs-keyword">var</span> = <span class="hljs-number">3</span> * <span class="hljs-keyword">var</span>;</code><ul style="" class="pre-numbering"><li>1</li><li>2</li><li>3</li><li>4</li><li>5</li><li>6</li><li>7</li></ul>

上述代碼在一個warp中執行的情況如下圖所示：

第一行所有線程執行int tid = threadIdx.x;。

所以第二行if(tid == 0) var = var + 1判斷,由于只有線程0執行var = var + 1;

第三行是線程0到線程31執行else var = var + 2;

第四行所有線程執行var = 3 * var;

上述代碼有兩個方面的缺陷：
1. cuda線程執行if,else語句的效率非常低。
2. 由于判斷產生了分支，導致第二行和第三行是串行執行的。
解決方法：
1. 通過查找表去掉分支
2. 通過計算去掉分支，例如：上述代碼可以轉換為：var = 3*(var+1+(var>0))

因此，我們在編程的時候要盡量是warp塊完美對齊，也就是說一個warp里都滿足條件或者都不滿足條件，如果實在無法對齊，也就是說產生分支的時候，可以用上述的方法解決分支問題。

總結

以上是生活随笔為你收集整理的CUDA分支优化的全部內容，希望文章能夠幫你解決所遇到的問題。

如果覺得生活随笔網站內容還不錯，歡迎將生活随笔推薦給好友。