Presja na rejestry i spillowanie

Rejestry to najszybsze miejsce przechowywania dostępne dla wątku, ale każdy SM ma ich ograniczoną liczbę. To, jak intensywnie kernel z nich korzysta, nazywa się register pressure.

Każdy aktywny wątek korzysta z jednego register file przypisanego do SM. Im więcej rejestrów potrzebuje każdy wątek, tym mniej wątków może jednocześnie pozostać aktywnych.

Duże wykorzystanie rejestrów bezpośrednio ogranicza liczbę warpów mieszczących się na SM. Ten spadek occupancy może pozostawić sprzętowi zbyt mało pracy do ukrywania opóźnień.

Gdy wątek potrzebuje więcej rejestrów, niż jest dostępnych, kompilator przenosi dodatkowe wartości poza nie. To przepełnienie jest określane jako register spill do wolniejszej pamięci.

Wartości wypchnięte z rejestrów trafiają do local memory, która znajduje się w pamięci DRAM poza układem. Każde takie przepełnienie zastępuje szybki dostęp do rejestru powolną podróżą tam i z powrotem.

Należy poprosić kompilator o zaraportowanie wykorzystania. Dodanie --ptxas-options=-v do nvcc wypisuje liczbę rejestrów na wątek oraz liczbę bajtów przepełnień dla każdego kernela.

Raport zawiera liczbę zapisów i odczytów przepełnień. Każda niezerowa wartość spill jest sygnałem ostrzegawczym, że kernel płaci za powolny ruch danych w pamięci lokalnej.

Limit można ustawić podczas kompilacji. Flaga --maxrregcount ogranicza liczbę rejestrów na wątek, poświęcając nieco szybkości na rzecz większej zajętości.

Wskazówka przypisana do konkretnego kernela jest często lepszym rozwiązaniem. __launch_bounds__ informuje kompilator o rozmiarze bloku, aby mógł zaplanować wykorzystanie rejestrów pod kątem oczekiwanej zajętości.

Często wystarczy przechowywać jednocześnie mniej wartości. Ponowne obliczenie taniego wyniku lub zawężenie scope zmniejsza liczbę aktywnych rejestrów.

ILP i rozwijanie pętli zwiększają presję, a limity zwiększają zajętość. Sztuka polega na znalezieniu balance, przy którym program działa najszybciej — i tylko profiler może wskazać, gdzie ona występuje.

Dowiedział się Pan/Pani, że duża register pressure zmniejsza zajętość i może powodować przepełnienia do wolnej pamięci DRAM. Należy mierzyć wykorzystanie, ograniczać liczbę rejestrów i korzystać ze wskazówek profilera. 🎯