CUDA分支优化

标签： cuda分支优化warpSM

2015-07-16 10:24 293人阅读 评论(0)收藏举报

分类：

CUDA（26）

版权声明：本文为博主原创文章，未经博主允许不得转载。

在CUDA中，分支会极大的减弱性能，因为ＳＭ没有分支预测，因此只能让束内线程在每个分支上都执行一遍，当然如果某个分支没有线程执行，就可以忽略，因此要减少分支的数目。可以简单的说：
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中执行的情况如下图所示：

1. 第一行所有线程执行int tid = threadIdx.x;。
2. 所以第二行if(tid == 0) var = var + 1判断,由于只有线程0执行var = var + 1;
3. 第三行是线程0到线程31执行else var = var + 2;
4. 第四行所有线程执行var = 3 * var;

上述代码有两个方面的缺陷：
1. cuda线程执行if,else语句的效率非常低。
2. 由于判断产生了分支，导致第二行和第三行是串行执行的。
解决方法：
1. 通过查找表去掉分支
2. 通过计算去掉分支，例如：上述代码可以转换为：var = 3*(var+1+(var>0))

因此，我们在编程的时候要尽量是warp块完美对齐，也就是说一个warp里都满足条件或者都不满足条件，如果实在无法对齐，也就是说产生分支的时候，可以用上述的方法解决分支问题。