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

lur*_*her 6 c++ optimization cuda opencl gpu-programming

大多数情况下,在CUDA或OpenCL程序中需要分支,例如:

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

代码总是(或至少在大多数情况下)以非分支样式重写:

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

权衡似乎是在单个warp槽中运行而不是在所有线程上进行更多计算(在第二种情况下,总和总是执行,有时值为零)

假设分支操作将为每个可能的分支采用多个warp插槽,可以预期第二个将始终优于第一个,现在我的问题是; 我是否可以依赖编译器来优化1)2)只要它有意义,或者没有广泛适用的标准,这意味着如果没有尝试和分析,通常无法确定哪一个更好?

小智 0

我对 CUDA 没有太多记忆,但你为什么不并行化你的循环呢?您应该使用原子操作[1]来添加计算。我希望这能帮到您!抱歉,如果情况并非如此。

  1. 原子操作:http://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/