c++ - CUDA:巨大的性能影响调用成员函数

Question

当我理解这个Robert Crovella's所以 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 代码时 - 我有一个性能非常好的大函数 [我用 #define 来“内联”代码].然后我根据 C++ 类的注释和文档对其进行了修改——鼓励——现在我的代码是 1'500 倍!较慢。

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

我可以改进什么或者是将代码“拆除”回糟糕的 C 代码的唯一解决方案？

当我将成员代码放入类声明时，代码并没有改变。这意味着，编译器已经“知道”成员函数的代码。并且，如果我更改优化级别，代码根本不会更改 -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 慢的情况。您的代码速度较慢只是因为它是在 Debug模式下显然编译的。

显然我怎么强调都不够。