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

Posted 2023-03-16

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

【英文标题】：OpenCL How efficient is the use of user-defined function in kernel code in terms of memory and performance

【发布时间】：2021-09-04 00:00:04

【问题描述】：

在 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 内存有任何影响，那么最好不要将它们编写为函数，而是在任何地方使用代码？

【问题讨论】：

你能用一个定义明确的例子来澄清你的问题吗？对我来说，当前的问题毫无意义。 “用户定义的函数”是什么意思。

@JérômeRichard 感谢您的评论！让我更新我的问题

Add 可能会在编译时内联（AFAIK 它是在 Nvidia 平台上完成的），因此在实践中可能没有开销。您可以使用 inline 关键字或在建议的答案中提供一些提示。内置指令不一定是低级指令，因此“内置指令”应该与“用户定义函数”一样快。但是，请记住，内置插件的实现通常是智能的，并且针对目标平台进行了精心优化。

@JérômeRichard 感谢您提供信息！那我就内联函数吧！

【参考方案1】：

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

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));

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

【讨论】：

再次感谢您回答我的许多问题！还有"function many times in a kernel"是什么意思是在代码中多次调用函数吗？

我的意思是，当函数大约 100 行或更长并且在内核中被调用大约 10 次或更频繁时，我不会内联。内联本身就像宏一样工作，将函数部分复制到调用函数的 括号中。但是内联比使用宏更安全；过度使用宏会做一些讨厌的事情。

再次非常感谢您！但是，如果我的内核以某种方式获得了 100% 的注册表使用率，那么内联会溢出到全局吗？

如果您一次使用太多私有变量，例如在私有内存空间中有一个 30x30 浮点矩阵，则寄存器可能会溢出到全局内存，从而降低性能。如果内核只是冗长，则重复使用相同的寄存器并且不会溢出。无论您是否内联，寄存器溢出都可能独立发生。

非常感谢！我关于这件事的所有问题现在都已经解决了！