CUDA堆栈帧的大小增加由__forceinline__

Question

当我与__forceinline__声明设备的功能，所述接头的输出这样的信息：

2> nvlink : info : Function properties for '_ZN3GPU4Flux4calcILj512EEEvv': 
2> nvlink : info : used 28 registers, 456 stack, 15776 bytes smem, 320 bytes cmem[0], 0 bytes lmem 

，没有它的输出是：

2> nvlink : info : Function properties for '_ZN3GPU4Flux4calcILj512EEEvv': 
2> nvlink : info : used 23 registers, 216 stack, 15776 bytes smem, 320 bytes cmem[0], 0 bytes lmem 

为什么的大小没有使用__forceinline__时，堆栈帧会变小吗？ 保持堆栈尽可能小是多么重要？ 谢谢你的帮助。

Answer 1

减少堆栈帧的主要原因是堆栈分配在驻留在片外设备内存中的本地内存中。这使得对堆栈的访问（如果没有被缓存）慢。

为了表明这一点，让我举个简单的例子。考虑这样的情况：

__device__ __noinline__ void func(float* d_a, float* test, int tid) { 
    d_a[tid]=test[tid]*d_a[tid]; 
} 

__global__ void kernel_function(float* d_a) { 
    float test[16]; 
    test[threadIdx.x] = threadIdx.x; 
    func(d_a,test,threadIdx.x); 
} 

注意，__device__函数声明__noinline__。在这种情况下，

ptxas : info : Function properties for _Z15kernel_functionPf 
    64 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas : info : Used 7 registers, 36 bytes cmem[0] 

即，我们有64字节的栈帧。相应的反汇编代码是

MOV R1, c[0x1][0x100]; 
ISUB R1, R1, 0x40; 
S2R R6, SR_TID.X;     R6 = ThreadIdx.x 
MOV R4, c[0x0][0x20]; 
IADD R5, R1, c[0x0][0x4]; 
I2F.F32.U32 R2, R6;     R2 = R6 (integer to float conversion)    
ISCADD R0, R6, R1, 0x2; 
STL [R0], R2;      stores R2 to test[ThreadIdx.x]         
CAL 0x50; 
EXIT ;        __device__ function part 
ISCADD R2, R6, R5, 0x2; 
ISCADD R3, R6, R4, 0x2; 
LD R2, [R2];       loads d_a[tid] 
LD R0, [R3];       loads test[tid] 
FMUL R0, R2, R0;      d_a[tid] = d_a[tid]*test[tid] 
ST [R3], R0;       store the new value of d_a[tid] to global memory 
RET ; 

正如你可以看到，test被存储并从全局存储器加载，形成堆栈帧（它是16 floats = 64 bytes）。

现在改变设备功能

__device__ __forceinline__ void func(float* d_a, float* test, int tid) { 
    d_a[tid]=test[tid]*d_a[tid]; 
} 

即，__device__功能改变从__noinline__到__forceinline__。在这种情况下，我们有

ptxas : info : Compiling entry function '_Z15kernel_functionPf' for 'sm_20' 
ptxas : info : Function properties for _Z15kernel_functionPf 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 

即，我们现在有一个空的堆栈帧。事实上，反汇编代码变为：

MOV R1, c[0x1][0x100];    
S2R R2, SR_TID.X;     R2 = ThreadIdx.x 
ISCADD R3, R2, c[0x0][0x20], 0x2;  
I2F.F32.U32 R2, R2;     R2 = R2 (integer to float conversion) 
LD R0, [R3];       R2 = d_a[ThreadIdx.x] (load from global memory) 
FMUL R0, R2, R0;      d_a[ThreadIdx.x] = d_a[ThreadIdx.x] * ThreadIdx.x 
ST [R3], R0;       stores the new value of d_a[ThreadIdx.x] to global memory 
EXIT ; 

正如你所看到的，迫使内联使编译器进行适当的优化，因此现在test完全从代码丢弃。

在上例中，__forceinline__的效果与您所遇到的相反，这也表明，没有任何进一步的信息，第一个问题就无法回答。