cuda：扭曲发散开销与额外算术

当然，扭曲发散，通过if and switch语句，在 GPU 上要不惜一切代价避免。

但是扭曲发散的开销是多少（仅调度some执行某些行的线程）与额外无用的算术？

考虑以下虚拟示例：

版本1：

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    int D = 0;

    if (A < 10)
        D = A*6;
    else if (A < 17)
        D = A*6 + B*2;
    else if (A < 26)
        D = A*6 + B*2 + C; 
    else 
        D = A*6 + B*2 + C*3;

    return D;
}

vs.

版本2：

__device__ int get_D (int A, int B, int C)
{
    //The value A is potentially different for every thread.

    return  A*6 + (A >= 10)*(B*2) + (A < 26)*C + (A >= 26)*(C*3);
}

我的真实场景更复杂（更多条件），但想法相同。

问题：

扭曲发散的开销（在调度中）是否如此之大以至于版本 1）比版本 2 慢？

版本 2 需要比版本 1 更多的 ALU，其中大部分都浪费在“乘以 0”上（只有少数条件计算结果为 1 而不是 0）。这是否会将有价值的 ALU 占用在无用的操作中，从而延迟其他扭曲中的指令？

通常很难提供此类问题的具体答案。影响2个案例对比分析的因素有很多：

• 你说 A 对于每个线程来说可能是不同的，但这种情况的真实程度实际上会影响比较。
• 总的来说，您的代码是否受计算限制或带宽限制肯定会影响答案。 （如果您的代码受带宽限制，则可能有no两种情况之间的性能差异）。
• 我知道你已经将 A、B、C 识别为整数，但看似无害的更改，例如使它们float可能会显着影响答案。

幸运的是，有一些分析工具可以帮助给出清晰、具体的答案（或者可能表明这两种情况之间没有太大区别）。您已经很好地识别了您关心的 2 个具体案例。为什么不对2进行基准测试？如果您想更深入地挖掘，分析工具可以提供有关指令重放（由于扭曲发散而产生）带宽/计算限制指标等的统计数据。

我不得不对这个笼统的声明表示异议：

当然，在 GPU 上要不惜一切代价避免通过 if 和 switch 语句发生扭曲发散。

这根本不是真的。机器处理发散控制流的能力实际上是feature这使我们能够使用更友好的语言（例如 C/C++）对其进行编程，并且实际上将其与其他一些不为程序员提供这种灵活性的加速技术区分开来。

与任何其他优化工作一样，您应该首先将注意力集中在繁重的工作上。您提供的这段代码是否构成了您的应用程序完成的大部分工作？在大多数情况下，将这种级别的分析工作投入到基本上是粘合代码或不属于应用程序主要工作的部分中是没有意义的。

如果这是您代码的大部分工作，那么分析工具确实是一种获得有意义的良好答案的强大方法，这些答案可能比尝试进行学术分析更有用。

现在我来回答你的问题：

扭曲发散的开销（在调度中）是否如此之大以至于版本 1）比版本 2 慢？

这将取决于实际发生的分支的具体水平。在最坏的情况下，如果有 32 个线程的完全独立的路径，机器将完全序列化，并且您实际上以峰值性能的 1/32 运行。线程的二元决策树类型细分不能产生这种最坏的情况，但肯定可以通过树的末尾来接近它。由于最后的线程完全发散，可能会观察到此代码的速度下降超过 50%，甚至可能下降 80% 或更高。但这在统计上取决于差异实际发生的频率（即它依赖于数据）。在最坏的情况下，我预计版本 2 会更快。

版本 2 需要比版本 1 更多的 ALU，其中大部分都浪费在“乘以 0”上（只有少数条件计算结果为 1 而不是 0）。这是否会将有价值的 ALU 占用在无用的操作中，从而延迟其他扭曲中的指令？

float vs. int可能实际上对这里有帮助，并且可能是您可以考虑探索的东西。但第二种情况（对我来说）似乎与第一种情况具有相同的比较，但有一些额外的乘法。在浮点情况下，机器每个时钟每个线程可以执行一次乘法，因此速度相当快。在 int 情况下，速度较慢，您可以根据架构查看具体的指令吞吐量here http://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html#arithmetic-instructions。我不会过度担心那种算术水平。再说一次，它可能会让完全没有区别如果您的应用程序受内存带宽限制。

另一种梳理所有这些的方法是编写内核来比较感兴趣的代码，编译为 ptx (nvcc -ptx ...) 并比较 ptx 指令。这可以更好地了解机器线程代码在每种情况下的样子，如果您只是执行指令计数之类的操作，您可能会发现两种情况之间没有太大区别（在这种情况下应该有利于选项 2） 。