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

Question

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

我正在编写一个应用程序来计算多体哈密顿量(量子系统的模拟)。计算高度依赖于条件表达式。

我之所以要减少这些语句，是因为它们会带来性能开销。 (整个 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 考虑根据 block 大小对内核进行模板化

这并不总是——事实上，并不经常——有用，但马克哈里斯在他的 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]);
 }

(假设检查返回一个 bool 值，或者失败时返回整数 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;
 }

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

slct

Select one source operand, based on the sign of the third operand.

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

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

3。通过争取 warp 范围内的共识来有效地避免分支

(另请参阅@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(); }
     ...
 }

这可确保所有经线(可能中间的经线除外)要么所有 channel 都满足条件，要么所有 channel 都不满足条件。这意味着在其他 channel 执行其他分支时，这些扭曲中的任何 channel 都将不得不闲置等待。

有关真实示例，请查看 Mark Harris 的 presentation 中的幻灯片 7-13。我在上面提到过。