【发布时间】：2016-05-10 11:42:45
【问题描述】：
当我正确理解 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 代码时 - 我有一个性能非常好的大函数[我用#define“内联”了代码]。然后我通过 cmets 和文档对它进行了调整 - 鼓励 - 使用类将其调整为 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
【问题讨论】：
-
你有多个文件中的代码吗？您是否使用单独的编译和链接/可重定位设备代码？
-
@RobertCrovella 类声明在
.h文件中，然后包含在.cu文件中，但整个类定义在 one 文件中。它是一步编译的，对于这个测试用例，我没有链接它。我只用参数调用了nvcc并运行了反汇编程序。因此，此版本的代码不涉及链接器。 -
我很遗憾，但这根本不可能。你做错了什么（看看反汇编，这看起来非常像调试模式）。请发布一个确实表现出这种行为的minimal reproducible example，我愿意吃掉我的键盘。顺便说一句，
inline关键字在这种情况下完全没有任何作用。 -
@HubertApplebaum 现在凌晨 2 点在欧洲。我明天将创建 MCVE。内联命令：是的，我知道；这从测试中“保留”，我现在将其删除。只是考虑一下：反汇编代码显示由 CUDA 生成的代码我也不敢相信。在发布此问题之前，我运行了许多不同的测试/星座。当我通过文件托管程序发布整个反汇编文件时会有帮助吗？
-
很高兴知道我今天不会吃我的键盘！ （不过我并不太担心）。
标签： c++ performance cuda member