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

Question

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

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

考虑以下虚拟示例：

版本 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;
}

对比

版本 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);
}

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

问题：

warp 分歧的开销（在调度中）是否如此之大以至于版本 1）比版本 2 慢？

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

Answer 1

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

• 您说每个线程的 A 可能不同，但真实程度实际上会影响比较。
• 总体而言，您的代码是受计算约束还是受带宽约束肯定会影响答案。 （如果您的代码受带宽限制，那么这两种情况可能没有性能差异）。
• 我知道您已将 A、B、C 识别为整数，但看似无害的更改（例如将它们设为 float）可能会显着影响答案。

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

我必须反对这个笼统的声明：

当然，在 GPU 上要不惜一切代价避免通过 if 和 switch 语句产生的 warp 分歧。

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

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

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

现在我来回答你的问题：

warp 分歧的开销（在调度中）是否如此之大以至于版本 1）比版本 2 慢？

这将取决于实际发生的特定分支级别。在最坏的情况下，对于 32 个线程具有完全独立的路径，机器将完全序列化并且您实际上以 1/32 的峰值性能运行。线程的二叉决策树类型细分不会产生这种最坏的情况，但肯定可以在树的末端接近它。由于最后的完全线程分歧，可能会观察到此代码超过 50% 的减速，可能是 80% 或更高的减速。但这在统计上取决于实际发生分歧的频率（即它取决于数据）。在最坏的情况下，我希望第 2 版更快。

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

float 与 int 在这里可能会有所帮助，也许你可以考虑探索一下。但是第二种情况（对我而言）似乎与第一种情况具有所有相同的比较，但有一些额外的乘法。在浮点情况下，机器可以在每个时钟的每个线程中进行一次乘法运算，因此速度非常快。在 int 情况下，它更慢，您可以看到具体的指令吞吐量取决于架构here。我不会过分担心那种程度的算术。同样，如果您的应用受内存带宽限制，根本没有区别。

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