处理CUDA中的大型switch语句

gam*_*erx 5 parallel-processing cuda switch-statement statements

我知道不建议在CUDA中进行分支,因为它会对性能产生负面影响.在我的工作中,我发现自己必须实现包含几十个案例的大型switch语句.

有没有人知道这会对性能造成多大影响.(官方文档不是很具体)也有人有更有效的方式来处理这部分吗?

Rog*_*ahl 6

GPU以32个为一组运行线程,称为warp.每当warp中的不同线程经过代码中的不同路径时,GPU必须多次运行整个warp,每个代码路径一次.

要处理这个称为warp divergence的问题,您需要排列线程,以便给定warp中的线程尽可能少地通过不同的代码路径.当你这样做的时候,你几乎必须咬紧牙关并接受因任何剩余的扭曲发散而导致的性能损失.在某些情况下,您可能无法执行任何操作来排列线程.如果是这样,并且如果不同的代码路径是内核或整体工作负载的重要部分,则该任务可能不适合GPU.

如何实现不同的代码路径并不重要.if-else,switch,预测(在PTX或SASS),分支表或其他任何东西-如果它归结为线程在不同的路径运行的扭曲,你对性能的打击.

每条路径有多少线程也无关紧要,只是经线中不同路径的总数.

以下是另一个有关详细信息的答案.


gee*_*eek 2

避免多次切换的一个好方法是实现函数表,并根据切换条件按索引从表中选择函数。CUDA 允许您在内核中的函数上使用函数指针__device__

  • 这如何帮助减少扭曲发散和相关的性能损失? (8认同)
  • 他的意思是编译器可以选择以多种方式实现 switch 语句。他不是在谈论编译器命令行选项。我不认为使用函数表的建议解决方案真的有帮助——你仍然需要分支。此外,您可以使用带有继承和虚函数的类层次结构来实现函数表——然后编译器会为您完成这一切。但这并不意味着它会更快。 (3认同)
  • 编译器有很多很多的实现选项,可以用来将 switch 语句转换为可执行代码。从相当于一系列 if/then/else 子句的机器,到具有范围检查的计算分支或查找表。或者多种策略的组合。因此,在任何人推测 CUDA 硬件将如何处理 switch 语句之前,我们需要更多地了解 switch 语句试图完成的任务。 (2认同)