Interpretowanie danych wyjściowych --ptxas-options = -v

Question

Próbuję zrozumieć użycie zasobów dla każdego z moich wątków CUDA dla ręcznie pisanego jądra.

Skompilowałem mój plik kernel.cu do pliku kernel.o z nvcc -arch=sm_20 -ptxas-options=-v

i mam następujący wynik

ptxas info : Compiling entry function '_Z12searchkernel6octreePidiPdS1_S1_' for 'sm_20' 
ptxas info : Function properties for _Z12searchkernel6octreePidiPdS1_S1_ 
    72 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads 
ptxas info : Used 46 registers, 176 bytes cmem[0], 16 bytes cmem[14] 

Patrząc na wyjściu powyżej, jest to prawdą, że

• każdy wątek CUDA używa 46 rejestrów?
• nie ma rejestru przechodzenia na lokalną pamięć?

Mam również pewne problemy ze zrozumieniem wyników.

• Moje jądro wywołuje wiele funkcji __device__. IS 72 bajty suma całkowita pamięci dla ramek stosów funkcji __global__ i __device__?

• Jaka jest różnica między 0 byte spill stores i 0 bytes spill loads

• Dlaczego informacji cmem (co zakładam jest stała pamięć) powtarzające się dwa razy z różnymi postaciami? W jądrze nie używam żadnej stałej pamięci . Czy to oznacza, że ​​kompilator jest, pod maską, czy powie firmie GPU, aby używał jakiejś stałej pamięci?

Answer 1

• Każda nić CUDA użyciu 46 rejestrów? Tak, poprawne
• Nie ma rejestru rozlewania się na pamięć lokalną? Tak, zgadza
• wynosi 72 bajtów suma-całkowite pamięci dla ramek stos funkcji __global__ i __device__? Tak, poprawne
• Jaka jest różnica między 0-bajtowymi sklepami rozlewowymi i 0 bajtami rozlanych ładunków?
  • Uczciwe pytanie, obciążenia mogą być większe niż sklepy, ponieważ można wylać wyliczoną wartość, załadować ją raz, odrzucić (to znaczy przechowywać coś innego w tym rejestrze), a następnie załadować ponownie (to znaczy ponownie użyć). Aktualizacja: Zauważ też, że liczba wyciek obciążenia/sklep jest na podstawie analizy statycznej, jak opisano przez @njuffa w komentarzach poniżej
• Dlaczego jest informacja dla CMEM (co zakładam jest stała pamięć) powtórzono dwukrotnie z różnymi liczbami? W jądrze nie używam żadnej stałej pamięci. Czy to oznacza, że ​​kompilator będzie pod maską informował GPU o używaniu stałej pamięci?
  • Stała pamięć służy do kilku celów, w tym __constant__ zmiennych i argumentów jądra, różne „banki” są stosowane, to zaczyna się trochę szczegółowe, ale tak długo jak używasz mniej niż 64KB dla __constant__ zmiennych i mniej niż 4KB dla argumentów jądra będziesz w porządku.