CUDA：不同 .cu 文件中的设备函数内联

Question

两个事实：CUDA 5.0允许您在不同的对象文件中编译CUDA代码以供稍后链接。CUDA架构2.x不再自动内联函数。

像往常一样，在C/C++中，我在functions.cu中实现了一个__device__ int foo()函数，并将其头文件放置在functions.hu中。其他CUDA源文件中调用函数foo。

当我检查functions.ptx时，我发现foo()会溢出到本地存储器。为了测试目的，我注释掉了foo()的所有内容，只是让它返回1; 根据.ptx仍然有东西溢出到本地存储器。（我无法想象它是什么，因为该函数什么也没做！）

但是，当我将foo()的实现移动到头文件functions.hu中并添加__forceinline__限定符时，就不会写入任何内容到本地存储器！

这里到底发生了什么？为什么CUDA不会自动内联这样一个简单的函数？

分离头文件和实现文件的整个重点是使我的生活更轻松地维护代码。但是，如果我必须将一堆函数（或所有函数）粘贴到头文件中并强制进行内联，那么这有点违背了CUDA 5.0不同编译单元的目的...

有没有什么方法可以解决这个问题？

简单的实际例子：

functions.cu:

__device__  int  foo
        (const uchar param0,
        const uchar *const param1,
        const unsigned short int param2,
        const unsigned short int param3,
        const uchar param4) 
{    
    return 1; //real code commented out.
} 

上述函数会溢出到本地内存。

functions.ptx：

.visible .func  (.param .b32 func_retval0) _Z45fooPKhth(
        .param .b32 _Z45foohPKhth_param_0,
        .param .b64 _Z45foohPKhth_param_1,
        .param .b32 _Z45foohPKhth_param_2,
        .param .b32 _Z45foohPKhth_param_3
)
{
        .local .align 8 .b8     __local_depot72[24];
        .reg .b64       %SP;
        .reg .b64       %SPL;
        .reg .s16       %rc<3>;
        .reg .s16       %rs<4>;
        .reg .s32       %r<2>;
        .reg .s64       %rd<2>;

- cmo

---

并非所有本地内存使用都表示溢出。调用的函数需要遵循ABI调用约定，其中包括创建位于本地内存中的堆栈帧。如果您使用编译器开关-Xptxas-v，则编译器会报告堆栈使用情况和溢出情况。我希望它能显示出有本地内存用于堆栈帧，但没有溢出。据我所知，目前无法在单独编译的对象文件边界上进行内联。 - njuffa

@njuffa 感谢您关于跨对象内联的评论。我对这方面并不了解。所以一般来说，将所有函数都实现在头文件中（并且__forceinline__所有函数）以保证内联是最佳实践吗？ - cmo

这取决于您的需求。对于具有完整代码库的长编译时间的大型项目，或者要构建真正的设备代码库，分离式编译非常有用。分离式编译和内联之间的权衡与主机代码相似（例如ABI、调用开销）。一些主机编译器可以在分离的编译单元之间进行内联，但目前CUDA中不存在这种功能。为了获得最佳性能，使用带有内联函数的头文件仍然是一个好方法，这就是CUDA 5.0中CUDA标准数学库的实现方式。 - njuffa

@njuffa 请随意在下面提交正式的“答案”... - cmo

编译器使用启发式算法来控制函数的内联。因此，即使它们在同一编译单元中，也不一定内联每个函数。您可以使用__forceinline__和__noinline__属性覆盖启发式算法。 - njuffa

Answer 1

并非所有本地内存使用都代表溢出。调用函数需要遵循ABI调用约定，其中包括创建一个在本地内存中的堆栈帧。当通过命令行开关 -Xptxas -v 传递nvcc编译器时，编译器会将堆栈使用和溢出作为其子组件报告。

目前（CUDA 5.0），CUDA工具链不支持跨编译单元边界的函数内联，就像一些主机编译器所做的那样。因此，在分离编译的灵活性（例如仅重新编译长时间编译的大型项目中的一小部分以及可能创建设备端库）与通常由于函数内联而产生的性能增益之间存在权衡（例如消除由于ABI调用约定而产生的开销，启用额外的优化，例如跨函数边界的常量传播）。

单个编译单元内的函数内联由编译器启发式控制，试图确定是否有望从性能上获得收益（如果可能的话）。这意味着并非所有函数都可以内联。程序员可以使用函数属性 __forcinline__ 和 __noinline__ 覆盖启发式。

- njuffa