OpenCL 在内核代码中使用用户定义函数在内存和性能方面的效率如何

问题描述

在 OpenCL C 内核代码中，默认内置函数很好，但用户定义函数呢？与内置的相比，它们是否有任何性能和内存下降？ 如果是这样的话， 在 __kernel void 中编写所述用户定义函数一次或多次更好吗？

例如：-

gentype clamp ( gentype x,gentype minval,gentype maxval)

以上是一个内置函数，对性能没有影响，也不会减少gpu l0/l1缓存

通过用户定义的函数，我的意思是这样的

int Add(int a,int b)
{
   return a + b;
}

这些函数是否对l0/l1内存有任何影响，如果有，那么最好不要将它们写成函数，而是到处使用代码？

解决方法

我通常内联所有函数，除非它们很长并且在内核中被多次调用。例如

float __attribute__((always_inline)) sq(const float x) {
    return x*x;
}

用于计算 x 的平方。内联函数不会为调用自身的函数带来额外的计算成本。但是，如果您在内核中多次内联一个很长的函数，程序集会爆炸并溢出到全局内存中，从而导致性能损失。在这种情况下，与函数本身的执行时间相比，函数调用造成的开销可以忽略不计。 最后，如果您没有明确地内联一个非常短的函数，在大多数情况下编译器会自动内联。与 for 函数相同，对于使用 #pragma unroll 展开的循环也是如此。

关于数学函数，大多数与硬件直接相关，只有少数例外。例如，计数前导零函数 int y = clz(x); 尽管被翻译成 clz PTX 指令，但没有专用硬件并且比用 int y = 31-(int)(as_uint((float)x)>>23); 模拟它慢。同理，虽然反平方根rsqrt(x)是在硬件中执行的，

float __attribute__((always_inline)) fast_rsqrt(const float x) {
    const uint i = 0x5F37642F-((*(uint*)&x)>>1);
    return *(float*)&i;
}

运行速度稍快，但精度较低。在大多数情况下，内置数学函数是最好的选择。