CUDA分支优化

在CUDA中，分支会极大的减弱性能，因为ＳＭ没有分支预测，因此只能让束内线程在每个分支上都执行一遍，当然如果某个分支没有线程执行，就可以忽略，因此要减少分支的数目。可以简单的说：
1. 同一个warp中的所有线程执行相同的命令。
2. 并不是所有线程都会执行。

产生分支的一个常见场景就是if,else语句的使用，比如常用的边界判定。下面从一个实力说起。看如下代码：

    int tid = threadIdx.x;    if(tid == 0)        var = var + 1;    else        var = var + 2;    var = 3 * var;

上述代码在一个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里都满足条件或者都不满足条件，如果实在无法对齐，也就是说产生分支的时候，可以用上述的方法解决分支问题。