cuda：设备函数内联和不同的.cu 文件

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

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

当我检查时functions.ptx， 我看到foo()溢出到本地内存。为了测试的目的，我评论了所有的内容foo()刚刚成功return 1;根据.ptx。 （我无法想象它是什么，因为该函数什么也不做！）

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

这里发生了什么？为什么CUDA不自动内联这么简单的函数？

单独的头文件和实现文件的全部目的是让我的代码维护更加轻松。但是如果我必须在标题中添加一堆函数（或全部）并且__forceinline__它们，那么它就违背了 CUDA 5.0 不同编译单元的目的......

有没有办法解决？

简单、真实的例子：

函数.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.
} 

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

函数.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 调用约定，其中包括在本地内存中创建堆栈帧。当 nvcc 传递命令行开关 -Xptxas -v 时，编译器将堆栈使用情况和溢出报告为其子组件。

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

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