4

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

像往常一样在 C/C++ 中,我已经实现了一个函数并将其头文件__device__ int foo()放在. 该函数在其他 CUDA 源文件中调用。functions.cufunctions.hufoo

当我检查时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>;
4

1 回答 1

4

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

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

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

于 2013-06-14T16:49:12.743 回答