CUDA C __device__ 函数中的 __forceinline__ 效果

Posted 2023-04-15

【中文标题】CUDA C __device__ 函数中的 __forceinline__ 效果

【英文标题】：__forceinline__ effect at CUDA C __device__ functions

【发布时间】：2013-11-22 18:00:03

【问题描述】：

对于在常规 C 编码中何时使用内联函数以及何时避免使用内联函数，有很多建议。 __forceinline__ 对 CUDA C __device__ 函数有什么影响？它们应该在哪里使用以及在哪里避免使用？

【参考方案1】：

通常nvcc 设备代码编译器会自行决定何时内联特定的__device__ 函数，一般来说，您可能不需要担心用__forceinline__ 装饰器/指令覆盖它。

cc 1.x 设备不具备与新设备相同的硬​​件功能，因此编译器通常会自动内联这些设备的函数。

我认为指定__forceinline__ 的原因与您可能已经了解主机C 代码的原因相同。当编译器可能不会以其他方式内联函数时（例如在 cc 2.x 或更新的设备上），它通常用于优化。如果您只调用一次函数，这种优化（即函数调用开销）可能可以忽略不计，但如果您在循环中调用该函数，确保它被内联可能会显着改善代码执行。

作为一个反例，内联和递归通常有禁忌症。对于调用自身的递归函数，我认为不可能处理任意递归和严格的内联。因此，如果您打算递归使用函数（在 cc 2.x 及更高版本中受支持），您可能不想指定 __forceinline__。

一般来说，我认为你应该让编译器为你管理这个。它将智能地决定是否内联函数。

【讨论】：

在表达式模板中，您希望通过使用更简单自然的数学语法来实现与手写代码相同的性能。对于这种情况，我需要使用 __forceinline__（和主机 __forceinline 对应）来保证这一点。