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__的效果與您所遇到的相反，這也表明，沒有任何進一步的信息，第一個問題就無法回答。