CUDA：巨大的性能影响调用成员函数

Question

当我理解这个 Robert Crovella's SO answer 正确时，GPU 编译器应该积极地内联函数 出于性能原因。

我这里有一个测试用例，它没有发生，甚至这个非常简单的函数也没有内联，每次调用成员函数时编译器都会生成：

__device__ auto foo::isMemberHighest( int iParameterBar ) -> bool
{
    return iParameterBar == 1;
}

运行 -cubin 参数并用 nvdiasm 反汇编它我得到这个输出：

//--------------------- .text._ZN27foo15isMemberHighestEi --------------------------
    .section    .text._ZN27foo15isMemberHighestEi,"ax",@progbits
    .sectioninfo    @"SHI_REGISTERS=7"
    .align  64
        .global         _ZN27foo15isMemberHighestEi
        .type           _ZN27foo15isMemberHighestEi,@function
        .size           _ZN27foo15isMemberHighestEi,(.L_969 - _ZN27foo15isMemberHighestEi)
_ZN27foo15isMemberHighestEi:
.text._ZN27foo15isMemberHighestEi:
        /*0000*/                   MOV R0, R6;
        /*0008*/                   MOV R5, R5;
        /*0010*/                   MOV R4, R4;
        /*0018*/                   MOV R4, R4;
        /*0020*/                   MOV R5, R5;
        /*0028*/                   MOV R4, R4;
        /*0030*/                   MOV R5, R5;
        /*0038*/                   MOV R0, R0;
        /*0040*/                   MOV R0, R0;
.L_605:
        /*0048*/                   ISUB R3.CC, R4, RZ;
        /*0050*/                   ISETP.NE.X.AND P0, PT, R5, RZ, PT;
        /*0058*/                   PSETP.AND.AND P0, PT, !P0, PT, PT;
        /*0060*/                   PSETP.AND.AND P0, PT, !P0, PT, PT;
        /*0068*/                   NOP;
        /*0070*/                   SSY `(.L_449);
        /*0078*/               @P0 BRA `(.L_450);
        /*0080*/                   BRA `(.L_450);
.L_450:
        /*0090*/                   NOP.S              (*"TARGET= .L_449 "*);
.L_449:
        /*0098*/                   ISETP.EQ.AND P0, PT, R0, 0x1, PT;
        /*00a0*/                   SEL R0, RZ, 0x1, !P0;
        /*00a8*/                   MOV R0, R0;
        /*00b0*/                   MOV R4, R0;
        /*00b8*/                   RET;
.L_606:
        /*00c0*/                   EXIT;
.L_604:
        /*00c8*/                   EXIT;
.L_451:
        /*00d0*/                   BRA `(.L_451);
.L_969:

在 /*0098*/ 和 /*00a0*/ 之间有比较命令，然后是 return。

我的 C++ 代码对该函数有 5 个成员调用，我在反汇编代码中看到恰好有 5 个对该函数的调用：

JCAL `(_ZN27foo15isMemberHighestEi);

我现在有这个问题：一开始 - 当我有一个纯 C 代码时 - 我有一个性能非常好的大函数 [我 "inlined" 代码 #define].然后我用 classes 对 C++ 的评论和文档进行了修改 - 鼓励 - 现在我的代码是 1'500！较慢。

之前 18m 次迭代需要大约 73ms - 现在 560k 次迭代需要 3'300ms！这意味着它慢了 1'500 倍，这自然非常令人沮丧。当然，这不是导致这种延迟的唯一一个成员函数。我有大约 10 个它们导致每次迭代有 50 个 call 语句 [包括函数开销]，显然这里是瓶颈。

我可以改进什么或者是 "dismantle" 代码返回糟糕的 C 代码的唯一解决方案？

当我将成员代码放入 class 声明时，代码没有改变。这意味着，编译器"knows"已经编译了成员函数的代码。并且，如果我更改优化级别 -O1 -O2 -O3!

，代码根本不会改变

更新：

使用此语句编译：

/usr/local/cuda-7.5/bin/nvcc -cubin -O3 -Xcompiler -Wall -Xcompiler -Wextra
   -Xcompiler -Werror -std=c++11 --compile --relocatable-device-code=false
   -gencode arch=compute_30,code=sm_30  -x cu -o CudaCore.cubin "../cuda/CudaCore.cu"
&& nvdisasm CudaCore.cubin > CudaCore.cubin.asm

Answer 1

将评论总结成某种答案：

我还没有看到 C++ 比 C 慢的情况。您的代码速度较慢只是因为它显然在调试模式下编译。

而且我不能再强调显然就够了。

CUDA：巨大的性能影响调用成员函数

CUDA: huge performance impact calling member functions

c++

performance

cuda

member