cuda inline 和 noinline 设备函数

Question

根据文档，在计算能力为 1.x 的设备中，编译器默认内联 __device__ 函数，但对于计算能力为 2.x 及更高版本的设备，只有在编译器认为合适的情况下才会这样做。什么时候不适合？还有限定符，例如__noinline__ 和__forceinline__。在哪些情况下最好不要内联 __device__ 函数？

Answer 1

用于内联的编译器启发式可能会评估内联的潜在性能优势，因为它消除了针对其他特征（包括编译时间）的函数调用开销。激进的内联会导致非常大的代码，从而导致非常长的编译时间。通过观察为许多不同内核生成的代码，CUDA 编译器似乎在绝大多数情况下是内联的。请注意，在某些情况下，当前无法进行内联，例如，当调用的函数位于不同的、单独编译的编译单元中时。

根据我的经验，覆盖编译器的内联启发式算法的情况很少见。我使用__noinline__ 来限制代码大小，从而减少过多的编译时间。我知道，使用__noinline__ 对注册压力没有可预测的影响。内联可能允许更激进的代码移动，例如加载调度，这可能会增加寄存器压力，而由于 ABI 对寄存器使用的限制，不内联可能会增加寄存器压力。我从来没有发现使用__noinline__ 可以提高性能的情况，但这种情况当然可能存在，可能是由于指令缓存效应。

Answer 2

我的经验是，如果你强制 __device__ 函数调用被内联编译，它可以将运行时间减少一半。就在最近的一次中，我进行了一个内联函数调用（仅将 5 个变量传递给函数），内核执行时间从 9.5 毫秒减少到 4.5 毫秒（几乎是一半）。如果你认为你想以一周或更长时间的总运行时间执行相同的内核数亿次（比如我的案例和许多其他从事 CFD 或 MD 项目的案例），与巨大的编译时间相比，编译时间的增加并不重要在运行时保存。

总而言之，我认为值得尝试内联函数调用对运行时的影响，尤其是对于运行时间很长的代码。