如何减少内核中的分支/if 语句的数量？

Question

有没有一种聪明的方法可以减少 CUDA 内核中 if statements 的数量？

我正在编写一个计算多体哈密顿量的应用程序（模拟量子系统）。计算高度依赖于conditional expressions。

我想减少这些语句的原因是它们引入了性能开销。 （整个warp 输入if(){} else if(){} 语句的每个选项 - 如果不满足条件，则thread 在给定时间内保持不活动状态。

问题：
1. switch() 语句能解决问题吗？
2. 下面的代码是为了代表一个总体思路：

class tag_option_1 {};
class tag_option_2 {};
class tag_option_3 {};

template<typename T> __device__
int cal_something(int ab, int cd)
{
    return -12345; // error value. default case is an error.
};

template<> __device__
int cal_something<tag_option_1>(int ab, int cd)
{
    // return something
}

template<> __device__
int cal_something<tag_option_2>(int ab, int cd)
{
    // return something
}

template<> __device__
int cal_something<tag_option_3>(int ab, int cd)
{
    // return something
}

////////////////////////////////
// version #1:

__global__
void calc_hamiltonian(int * foo, int * bar)
{
    unsigned int tid = /* calce thread index*/;

    // do something... 

    if (/* condition */)
    {
        cal_something<tag_option_1>(foo[tid], bar[tid]);
    }
    else if(/* condition */)
    {
        cal_something<tag_option_2>(foo[tid], bar[tid]);
    }
    else if(/* condition */)
    {
        cal_something<tag_option_3>(foo[tid], bar[tid]);
    }
    // no default case.

    // do something... 
}

////////////////////////////////
// version #2:

// "magical" way to select a version:
// variant is meant to be something like "boost::variant" but implementented without a single "if" statement.
// This "magical" step is meant to be resolved at compile time.
__devcie__
variant <tag_option_1, tag_option_2, tag_option_3> 
version_selector(int ab, int cd)
{
    // magic happens here.
}


__global__
void calc_hamiltonian(int * foo, int * bar)
{
    unsigned int tid = /* calce thread index*/;

    // do something... 

    cal_something <version_selector(foo[tid], bar[tid])> (foo[tid], bar[tid]);

    // do something... 
}

有没有办法实现上面例子的version #2，或者在CUDA C/C++中是不可能的？

Answer 1

我一般同意@njuffa 的建议，即不要试图人为地扭曲您的自然编码风格，并且您应该追求性能（以及可读性和可维护性）而不是计算源代码中的分支。特别是因为编译器有时可能会让它们消失。

话说回来……

一些常见的“聪明”方法来减少分支的数量（在 CUDA 和一般情况下）：

1。提前安排条件检查，最好在编译时进行

举例说明。版本 1：

 void foo(int* a, bool cond) {
     ...
     for(int i = 0; i < lots; i++) {
         if (cond) do_stuff()
         else do_other_stuff();
     }
     ...
 }

 bool cond = check_stuff();
 foo(data, cond);

版本 2：

 void foo(int* a, bool cond) {
     ...
     if (cond) {
         for(int i = 0; i < lots; i++) { do_stuff(); }
     }
     else {
         for(int i = 0; i < lots; i++) { do_other_stuff(); }
     }
     ...
 }

 bool cond = check_stuff();
 foo(data, cond);

版本 3：

 template <bool Cond>
 void foo(int* a) {
     ...
     if (cond) {
         for(int i = 0; i < lots; i++) { do_stuff(); }
     }
     else {
         for(int i = 0; i < lots; i++) { do_other_stuff(); }
     }
     ...
 }

 bool cond = check_stuff();
 if (cond) foo<true>(data) else foo<false>(data);

版本 4：

 template <bool Cond>
 void foo(int* a) {
     ...
     for(int i = 0; i < lots; i++) {
         if (cond) do_stuff()
         else do_other_stuff();
     }
     ...
 }

 bool cond = check_stuff();
 if (cond) foo<true>(data) else foo<false>(data);

第 3 版和第 4 版的好处在于，虽然它们看起来有一个分支，但实际上并没有——编译器要么只接受“then”语句，要么只接受"else" 语句，但不能同时在同一个函数中。

从版本 1 到版本 2 是编译器可能会为您做的事情；但有时它并不像示例中那样简单，您必须自己处理。从版本 2 到版本 3 是编译器永远不会为您做的事情。转到第 4 版有点像绕了一圈，因为它看起来像第 1 版，没有代码重复 - 但分支实际上仍然消失了。

1.1 考虑根据块大小对内核进行模板化

这并不总是——事实上，不是经常——有用，但是 Mark Harris 在他的presentation 中给出了一个关于使用 CUDA 优化并行缩减的众所周知的例子。查看幻灯片 24-27 中的优化 #6。但是不要尝试这样的事情 - 这很难看而且有些脆弱 - 除非你已经仔细计时你的执行以确保它是值得的。

2。使行为在数据中而不是在控制流中发散

版本 1：

 void foo(int* a, int *b) {
     ...
     if (check(a[global_thread_index]) { b[global_thread_index]++; }
 }

版本 2：

 void foo(int* a, int *b) {
     ...
     b[global_thread_index] += check(a[global_thread_index]);
 }

（假设 check 返回一个布尔值，或者一个整数 0 表示失败，1 表示成功。）

这里我不太确定 CUDA 编译器会做什么；另外，您通过编写此代码并可能破坏"principle of least astonishment" 来支付可读性损失。但是你可以找到不那么做作的例子。

2.1 将分支限制为值的三元运算选择

还有一个版本 3：

 void foo(int* a, int *b) {
     ...
     b[global_thread_index] = check(a[global_thread_index]) ? 1 : 0;
 }

现在，这仍然有一个分支 - 三元运算符只是“如果”的简写，但如果你可以让你的代码达到这种状态，分歧将仅限于单个语句和每个分支，甚至可能更少，如果 CUDA 编译器设法使用 slct PTX statement:

slct

根据第三个操作数的符号选择一个源操作数。

这将分支的语义“包装”到单个指令的组合逻辑中。

当然，slct 可能会在其他情况下被编译器使用；这不取决于你。

3。通过争取全域共识来有效避免分支

（另请参阅@RobertCrovella 的相同效果的评论。）

再次，举例说明。

 void foo(int* a, int *b) {
     ...
     if (threadIdx.x % 2 == 0) { do_stuff(); }
     else { do_other_stuff(); }
     ...
 }

版本 2：

 void foo(int* a, int *b) {
     ...
     if (threadIdx.x >= blockDim.x / 2) { do_stuff(); }
     else { do_other_stuff(); }
     ...
 }

这可确保所有经线（可能除了中间的经线）的所有通道都满足条件或所有通道不满足条件。这意味着这些经线中的任何通道都不会在其他通道执行另一个分支时必须空闲等待。

对于现实生活中的示例，请查看上面提到的 Mark Harris 的presentation 中的幻灯片 7-13。