cuda：设备功能内联和不同的 .cu 文件答案

【问题标题】：cuda: device function inlining and different .cu filescuda：设备功能内联和不同的 .cu 文件
【发布时间】：2013-06-13 18:14:44
【问题描述】：

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

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

当我检查functions.ptx 时，我看到foo() 溢出到本地内存。出于测试目的，我评论了foo() 的所有内容，然后将其设为return 1;，根据.ptx，某些内容仍会溢出到本地内存。（我无法想象它是什么，因为这个函数什么都不做！）

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

这是怎么回事？ 为什么 CUDA 不自动内联这么简单的函数？

单独的头文件和实现文件的全部意义在于让我更轻松地维护代码。但是，如果我必须在标头中添加一堆函数（或所有函数）并 __forceinline__ 它们，那么这有点违背了 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>;

【问题讨论】：

并非所有本地内存使用都表示溢出。被调用的函数需要遵循 ABI 调用约定，其中包括在本地内存中创建堆栈帧。如果您使用编译器开关 -Xptxas -v，编译器会报告堆栈使用情况和溢出。我希望这表明有用于堆栈帧的本地记忆库但没有溢出。据我所知，内联目前不能跨越单独编译的目标文件的边界。
@njuffa 你说对了一部分。像我提到的那样，还有更多麻烦的功能。正如您所建议的，其中一些使用堆栈帧但不会溢出：24 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 但是，其他函数确实从本地内存调用：24 bytes stack frame, 24 bytes spill stores, 24 bytes spill loads
@njuffa 谢谢关于跨对象内联的评论。我对那东西了解不多。所以一般来说，最好的做法是在头文件中实现我的所有函数（以及__forceinline__所有这些函数）以保证内联？
取决于您的需求。对于完整代码库的编译时间较长的大型项目或构建真正的设备代码库，单独编译非常有用。单独编译和内联之间的权衡类似于它们对主机代码的权衡（例如 ABI、调用开销）。一些主机编译器提供跨单独编译的编译单元的内联，但目前 CUDA 中不存在该功能。为了获得最佳性能，使用带有内联函数的头文件仍然是一个好方法，这就是 CUDA 标准数学库在 CUDA 5.0 中的实现方式。
@njuffa 所以如果我将所有这些 CUDA 函数都放在头文件中，那么就没有必要为它们使用__forceinline__ 限定符，对吧？

标签： cuda gpu inline nvidia ptx

【解决方案1】：

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

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

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

【讨论】：