3

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

但是，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);
}

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

问题：

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

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

4

1 回答 1

3

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

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

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

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

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

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

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

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

现在我来回答你的问题：

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

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

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

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

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

于 2013-05-24T16:38:00.023 回答