CUDA /openCL;将分支重写为非分支表达式

CUDA /openCL; rewriting branches as non-branching expression

本文关键字：分支 重写 表达式 openCL CUDA      更新时间：2023-10-16

大多数情况下，CUDA 或 OpenCL 程序中需要一个分支，例如：

for (int i=0; i<width; i++)
{
   if( i % threadIdx.x == 0)
     quantity += i*i;
}

代码总是可以(或至少在大多数时候(以非分支样式重写：

for (int i=0; i<width; i++)
{
   quantity += i*i* (i % threadIdx.x != 0);
}

权衡似乎是在单个 warp 插槽中运行，还是在所有线程上执行更多计算(在第二种情况下，总和始终执行，只是有时值为零(

假设分支操作将为每个可能的分支占用多个经线槽，人们会期望第二个始终比第一个更好，现在我的问题是;我可以依靠编译器将 1( 优化为 2( 只要有意义，或者没有广泛适用的标准，这意味着如果不尝试和分析，就无法确定哪个更好？

模运算相当昂贵：我有理由确信添加模比只执行 1 个线程的单个指令会花费更多的时间。您的单个分支语句(一个没有else的if(只会在执行 if 语句时挂起其他线程。由于 GPU 针对非常快速的上下文切换进行了优化，因此成本应该很小。

但是，建议您不要使用长分支语句：GPU上的串行计算过多(即一个线程完成所有工作(会抵消并行性的优势。

根据我的经验 - 完全取决于编译器编写者来优化这些边缘情况。

那么我能想到任何 1( 不能变成 2( 的情况吗？这里有一个：我编写了内核，其中每 10 个线程运行某些计算部分或类似的东西更有效，在这种情况下，即使存在数学运算(除法减法(也无法推断出这种优化可以产生相同的结果，无论条件如何与"在所有线程上运行但产生零结果"。

但是，即使考虑到检查 threadId == 0 是一种足够常见的情况，我也不知道它是否真的进行了优化。我敢打赌，这取决于实现甚至设备本身(CPU 与 GPU(。

您将不得不尝试它才能真正找出最有效的方法，不仅因为上述原因，还因为工作计划程序的行为可能会有所不同，具体取决于调度/启动/停止一组线程的成本，而不是让它们全部运行(并且大多数提供零/标识结果(。

希望这有帮助！

我对 CUDA 没有太多记忆，但你为什么不并行化你的循环？你应该使用原子操作[1]来添加你的计算。希望对您有所帮助！对不起，如果不是这样。

  1. 原子操作：http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/
相关文章：