6

大多数情况下，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);
}

权衡似乎是在单个扭曲槽中运行而不是在所有线程上进行更多计算（在第二种情况下，总和总是执行，只是有时值为零）

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

4

3 回答 3

3

模运算相当昂贵：我有理由确定添加模会比仅使用只有 1 个线程执行的单个指令花费更多时间。您的单个分支语句（if带有 no else）只会在执行 if 语句时挂起其他线程。因为 gpus 针对非常快速的上下文切换进行了优化，所以成本应该很低。

但建议您不要使用长分支语句：GPU 上过多的串行计算（即一个线程完成所有工作）否定了并行性的优势。

于 2012-05-15T21:59:13.433 回答
1

以我的经验 - 完全由编译器编写者来优化这些边缘情况。

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

但是，即使检查 threadId == 0 是一个足够常见的场景，我也不知道它是否真的针对它进行了优化。我敢打赌，这取决于实现，甚至取决于设备本身（CPU 与 GPU）。

您将不得不尝试它才能真正找出最有效的方法，这不仅是因为上述原因，还因为工作调度程序可能会根据调度/启动/停止一组线程的成本而不是拥有它们都运行（并且大多数提供零/身份结果）。

希望这可以帮助！

于 2012-05-15T21:21:22.630 回答
0

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

  1. 原子操作：http ://supercomputingblog.com/cuda/cuda-tutorial-4-atomic-operations/
于 2012-05-15T20:55:41.353 回答