Ograniczenia rejestrów i pamięci współdzielonej

Każdy SM ma stałą pulę rejestrów i pamięci współdzielonej. Każdy aktywny blok musi wydzielić swoją część z tych pul.

Kernel wykorzystuje określoną liczbę rejestrów na wątek. Po pomnożeniu jej przez liczbę wątków w bloku otrzymuje się koszt rejestrowy jednego bloku.

Jeśli każdy wątek potrzebuje wielu rejestrów, mieści się mniej wątków, a więc mniej bloków może pozostać aktywnych. Kernele intensywnie korzystające z rejestrów mają niższe occupancy.

Po skompilowaniu z tą flagą nvcc wyświetla liczbę rejestrów na wątek, co pozwala zauważyć, że kernel jest zbyt żądny rejestrów.

Można wymusić limit za pomocą launch bound, aby kompilator używał mniej rejestrów i pozwolił większej liczbie warpów pozostać aktywnymi.

Każdy blok może zażądać pamięci współdzielonej. SM mieści tylko tyle bloków, na ile pozwala jego pula pamięci współdzielonej, często od 48 do 100 KB.

Zażądanie dużego kafelka __shared__ sprawia, że na jednym SM mieszczą się tylko jeden lub dwa bloki. Duże kafelki wymieniają occupancy na ponowne wykorzystanie danych.

SM oblicza liczbę bloków dozwoloną przez rejestry, pamięć współdzieloną i limit warpów. O occupancy decyduje najmniejsza z tych wartości.

Gdy wątek potrzebuje więcej rejestrów, niż mu przydzielono, nadmiar jest przenoszony do wolnej pamięci lokalnej. Takie przenoszenie może zaszkodzić bardziej niż niskie occupancy.

Zmniejszenie liczby rejestrów lub pamięci współdzielonej zwiększa occupancy, ale zbyt agresywne ograniczenie powoduje przenoszenie rejestrów. Najlepszy punkt wymaga zachowania równowagi.

Przed optymalizacją należy zawsze odczytać rzeczywiste użycie rejestrów i pamięci współdzielonej z danych kompilatora. Zgadywanie zwykle prowadzi do wyboru niewłaściwego parametru.

Przypomnijmy sobie, jak SM określa liczbę aktywnych bloków.

Widać już, że rejestry i pamięć współdzielona są stałymi budżetami SM, a najciaśniejsze ograniczenie wyznacza maksymalne occupancy. Należy uważać na przenoszenie rejestrów. 🧮