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

OpenCL How efficient is the use of user-defined function in kernel code in terms of memory and performance

在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的平方。内联函数不会为调用自身的函数带来额外的计算成本。但是,如果您在内核中多次内联一个非常长的函数,程序集就会崩溃并溢出到全局内存中,从而导致性能下降。在这种情况下,与函数本身的执行时间相比,函数调用的开销可以忽略不计。 最后,如果你没有明确地内联一个非常短的函数,编译器在大多数情况下会自动完成。对于使用 #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) {
    return as_float(0x5F37642F-(as_int(x)>>1));
}

运行速度稍快但不太准确。在大多数情况下,内置数学函数是最好的选择。