c++ - CUDA/开放式；将分支重写为非分支表达式

Question

大多数时候在 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)，或者没有广泛适用的标准，这意味着如果不进行尝试和分析就无法总体上决定哪个更好？

Answer 1

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

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