NVIDIA GPU 支持分支预测？（使用 OpenACC）答案

【问题标题】：NVIDIA GPU support branch prediction? (with OpenACC)NVIDIA GPU 支持分支预测？（使用 OpenACC）
【发布时间】：2015-12-12 11:22:44
【问题描述】：

我正在使用带有 OpenACC 的 NVIDIA GPU （NVIDIA GeForce960，编译器：PGI 15.7）

NVIDIA GPU 是否支持分支预测？我的代码在长循环中有条件执行代码。但是当我在 GPU 上运行我的代码时，它需要很长时间。

下面是示例代码。

#pragma acc parallel loop
for(i=0; i<1000; i++)  // NVIDIA GPU kernel with OpenACC
   for(j=0; j<1000; j++)
      if(a[i][j]==value)
         // do something...
      else if(a[i][j]==value2)
         // do another something...
      else
         // do another something...

如果此代码在 CPU 上编译并运行，CPU 将对条件执行循环进行分支预测。 NVIDIA GPU 是否具有与 CPU 的分支预测相同或相似的技能？

如果有，我该如何启用它？

补充问题：

PGI 15.7 编译器是否可以为 GPU 编译优化代码？我知道通常的编译器（gcc...等）可以使用优化技能进行编译，例如更改条件代码执行顺序。

【问题讨论】：

当前 CUDA gpus 不支持分支预测。可能不是潜水到这么低的水平，而是你关心的事情是“为什么我的代码运行缓慢？”但是没有人可以根据您发布的内容来回答这个问题。
感谢您的回复，Robert :) 我真正想知道的是“CUDA gpu 可以做分支预测吗？”。如果可以，是否有一些开关可以启用它。但正如你所回答的，现在我知道 CUDA gpu 没有分支预测。

标签： cuda nvidia openacc

【解决方案1】：

NVIDIA GPU 与 CPU 的分支预测有相同或相似的技能吗？

当前的 NVIDIA GPU 不支持分支预测。

PGI 15.7 编译器是否可以为 GPU 编译优化代码？

是的，PGI 工具可以执行各种优化。这通过-Ox 命令行开关（例如-O3）在高级别的控制，就像gcc/g++ 一样。这种优化可以发生在编译的各个阶段，例如将 OpenACC 源代码转换为 CUDA PTX，以及将 CUDA PTX 转换为 CUDA SASS。

【讨论】：

感谢 Robert 的回复 :) 与 -O3 选项相比，PGI 编译器可以重新排序指令以优化条件执行区域吗？（即 if-else）OpenACC 规范没有像 CUDA C 那样指示块和线程索引的特定方法，我无法在 1 个扭曲中消除分支分歧。我认为它发生了很多“非活动”状态线程，并减慢了我的代码。
仅供参考：PGI 编译器的推荐编译器优化标志是 -fast，其中包括 -O2 和许多其他有益的优化。您也可以将其与-O3 结合使用，但我认为我从未见过-O3 -fast 的表现优于-fast.
有什么方法可以重构你的代码，让分支语句发生在最内层循环之外？它通常是在 OpenACC 中被矢量化的最内层循环，因此这是您要避免分歧的地方。如果不能完全去除，有时可以重构循环，让所有真假迭代都归为一组，这样实际上只有一个warp是发散的。
感谢 jeff :) 不幸的是，我无法修复我的整个代码。因为这是我的同事编写的用于图像处理的非常长且复杂的代码（而且我无法理解所有内容）。所以对于修复代码，我必须知道哪些循环会产生许多分支分歧。你知道一些工具可以找到这些东西吗？
nvidia 分析工具应该涵盖这一点。 NSight/Visual 分析器应该在 GUI 中显示它。 nvprof 具有指标warp_execution_efficiency 和warp_nonpred_execution_efficiency。