我应该创建多个 OpenCL 内核来避免条件语句吗？答案

【问题标题】：Should I create multiple OpenCL kernels to avoid conditional statements?我应该创建多个 OpenCL 内核来避免条件语句吗？
【发布时间】：2011-05-05 18:57:54
【问题描述】：

在 OpenCL 中，我有一个需要对复杂和真实数据进行操作的内核。我可以在其中放入一个条件语句，调用正确的代码行来处理这个问题，或者我可以有两个我调用的内核并将条件语句推送到我的调用代码中。

这显然不利于可维护性，但对性能有影响吗？

【问题讨论】：

标签： performance cuda opencl nvidia conditional-statements

【解决方案1】：

如果它只是一个条件语句，根据我的经验，性能差异绝对可以忽略不计，至少在 NVidia 硬件上是这样。

基本上，只要所有（或大多数）工作项遵循相同的代码路径，就可以了。由于在您的情况下采用的代码路径取决于内核参数，因此所有工作项都遵循相同的路径。

【讨论】：

【解决方案2】：

稍微取决于条件的位置。首先是代码的可读性，然后是性能，在你测量它并发现它是一个问题之后

例如。 kernel_for_RGB_image 和 kernel_for_ABGR_image 似乎是一个合理的使用，不同的内核有效地展开一些深层内部循环可能是一个更大的维护难题。

【讨论】：

【解决方案3】：

我认为最好的方法是实际尝试和基准测试两个变体。在某些情况下，编译了多个条件块，即使只执行其中一个，也会导致性能下降。原因是 GPR（通用寄存器）：编译器分配尽可能多的寄存器，以满足最坏情况的需要。

我可以建议这样的解决方案：有一个内核函数，但有编译时条件：

__kernel void work()
{
#if VAR
    // one code
#else
    // another code
#endif
}

然后你需要在改变条件时将true/false设置为VAR重新编译内核。显然，对于编译器而言，它与两个内核没有什么不同，但如果这些内核的一部分代码相同，那么维护可能会更好。

【讨论】：