NVCC寄存器使用情況報告

Question

我試圖讓使用NVCC選項
--ptxas-options=v我CUDA內核有關的寄存器使用情況的信息，並同時與全球功能一切正常，我有一些困難由於

ptxas info : Used N registers

線設備那些缺少在輸出中。我試圖使用noinline關鍵字，並將它們保存在另一個文件中，與調用全局函數有關，因爲我認爲NVCC報告了全局函數的全部註冊使用情況，包括內聯後的被調用設備的全局函數，但沒有任何變化。我可以獲得關於設備功能的寄存器使用情況的信息，只將它們定義爲全局的。

你有什麼建議嗎？

謝謝！

Answer 1

據我所知，ptxas（設備彙編器）只輸出一個寄存器計數，它鏈接的代碼。獨立的__device__函數沒有被彙編器鏈接，它們只被編譯。因此，彙編器不會爲器件功能發出寄存器計數值。我不相信這有一個解決方法。

但是，仍然可以通過使用cuobjdump從彙編程序輸出中轉儲精靈數據來獲得__device__函數的寄存器佔用空間。你可以這樣做如下：

$ cat vdot.cu 
__device__ __noinline__ float vdot(float v1, float v2) { 
    return (v1 * v2); 
} 

__device__ __noinline__ float vdot(float2 v1, float2 v2) { 
    return (v1.x * v2.x) + (v1.y * v2.y); 
} 

__device__ __noinline__ float vdot(float4 v1, float4 v2) { 
    return (v1.x * v2.x) + (v1.y * v2.y) + (v1.z * v2.z) + (v1.w * v2.w); 
} 

$ nvcc -std=c++11 -arch=sm_52 -dc -Xptxas="-v" vdot.cu 
ptxas info : 0 bytes gmem 
ptxas info : Function properties for cudaDeviceGetAttribute 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for _Z4vdotff 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessor 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for _Z4vdot6float4S_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for cudaMalloc 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for cudaGetDevice 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for _Z4vdot6float2S_ 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for cudaOccupancyMaxActiveBlocksPerMultiprocessorWithFlags 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Function properties for cudaFuncGetAttributes 
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 

在這裏，我們有一個設備對象文件單獨編譯組三個__device__功能。運行在它cuobjdump會排出大量的輸出，但它，你將獲得每個功能的寄存器計數：

$ cuobjdump -elf ./vdot.o 

Fatbin elf code: 
================ 
arch = sm_52 
code version = [1,7] 
producer = cuda 
host = linux 
compile_size = 64bit 
compressed 

<---Snipped---> 


.text._Z4vdotff 
bar = 0 reg = 6 lmem=0 smem=0 
0xfec007f1 0x001fc000 0x00570003 0x5c980780 
0x00470000 0x5c980780 0x00370004 0x5c680000 
0xffe007ff 0x001f8000 0x0007000f 0xe3200000 
0xff87000f 0xe2400fff 0x00070f00 0x50b00000 

在輸出設備功能的第二行dot(float, float)可以看到函數使用6寄存器。這是我知道檢查器件功能寄存器佔位情況的唯一方法。