NVCC Registerverwendungsbericht in __device__ Funktion

Question

Ich versuche, einige Informationen über Registerverwendung in meinem CUDA Kernel NVCC mit Option
--ptxas-options=v zu bekommen und während sie mit globalen Funktionen alles in Ordnung ist, ich habe einige Schwierigkeiten mit mit das Gerät diejenigen der seit

ptxas info : Used N registers

Leitung wird in der Ausgabe fehlt. Ich habe versucht, das Noinline-Schlüsselwort zu verwenden und sie in einer anderen Datei in Bezug auf die aufrufende globale Funktion zu behalten, da ich dachte, dass NVCC die vollständige Registerbenutzung der globalen Funktion einschließlich der aufgerufenen Geräte nach dem Inline aber nichts gemeldet hat Änderungen. Ich kann die Informationen über die Registerbenutzung der Gerätefunktionen nur als global definieren.

Haben Sie Vorschläge?

Danke!

Answer 1

Wie ich es verstehe, gibt ptxas (der Gerät Assembler) nur eine Registeranzahl auf Code, der es verbindet. Standalone __device__ Funktionen sind nicht vom Assembler verknüpft, sie werden nur kompiliert. Daher gibt der Assembler keinen Zählwert für Gerätefunktionen aus. Ich glaube nicht, dass es einen Workaround dafür gibt.

Es ist jedoch immer noch möglich, den Registrierungs-Footprint einer __device__-Funktion zu erhalten, indem Sie die Elf-Daten aus der Assembler-Ausgabe mit cuobjdump ausgeben. Sie können dies wie folgt:

$ 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 

Hier haben wir einen separat kompilierten Satz von drei __device__ Funktionen in einem Gerät Objektdatei haben. Laufen cuobjdump auf es wird eine Menge von Ausgang emittieren, aber in es werden Sie einen Registerzählwert für jede Funktion erhalten:

$ 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 

In der zweiten Zeile der Ausgabe für die Gerätefunktion dot(float, float) Sie die Funktion 6 sehen können verwendet Register. Dies ist die einzige Möglichkeit, die ich kenne, um Fußabdrücke von Gerätefunktionsregistern zu untersuchen.