gam*_*erx 5 parallel-processing cuda switch-statement statements
我知道不建议在CUDA中进行分支,因为它会对性能产生负面影响.在我的工作中,我发现自己必须实现包含几十个案例的大型switch语句.
有没有人知道这会对性能造成多大影响.(官方文档不是很具体)也有人有更有效的方式来处理这部分吗?
GPU以32个为一组运行线程,称为warp.每当warp中的不同线程经过代码中的不同路径时,GPU必须多次运行整个warp,每个代码路径一次.
要处理这个称为warp divergence的问题,您需要排列线程,以便给定warp中的线程尽可能少地通过不同的代码路径.当你这样做的时候,你几乎必须咬紧牙关并接受因任何剩余的扭曲发散而导致的性能损失.在某些情况下,您可能无法执行任何操作来排列线程.如果是这样,并且如果不同的代码路径是内核或整体工作负载的重要部分,则该任务可能不适合GPU.
如何实现不同的代码路径并不重要.if-else
,switch
,预测(在PTX或SASS),分支表或其他任何东西-如果它归结为线程在不同的路径运行的扭曲,你对性能的打击.
每条路径有多少线程也无关紧要,只是经线中不同路径的总数.
避免多次切换的一个好方法是实现函数表,并根据切换条件按索引从表中选择函数。CUDA 允许您在内核中的函数上使用函数指针__device__
。