Słaba wydajność podczas wywoływania cudaMalloc z dwoma procesorami graficznymi jednocześnie

Question

Mam aplikację, w której dzielę obciążenie przetwarzania pomiędzy procesory GPU w systemie użytkownika. Zasadniczo istnieje wątek procesora na GPU, który inicjuje interwał przetwarzania GPU po uruchomieniu okresowym przez główny wątek aplikacji.

Weź pod uwagę następujące zdjęcie (wygenerowane przy użyciu narzędzia profilującego CUDA NVIDIA) na przykładzie interwału przetwarzania GPU - tutaj aplikacja korzysta z pojedynczego procesora graficznego.

Jak widać duża część czasu przetwarzania GPU jest zużywana przez dwie operacje sortowania i używam biblioteki Thrust dla tego (oporowej :: sort_by_key). Poza tym wygląda na to, że ciąg :: sort_by_key wywołuje kilka cudaMallocs pod maską, zanim rozpocznie właściwy sort.

Rozważmy teraz ten sam przedział przetwarzanie gdzie aplikacja rozłożone obciążenie przetwarzania ponad dwa GPU:

W idealnym świecie można oczekiwać odstęp przetwarzania 2 GPU być dokładnie połowę z pojedyncze GPU (ponieważ każdy GPU wykonuje połowę pracy). Jak widać, nie dzieje się tak częściowo dlatego, że cudaMalloki wydają się trwać dłużej, gdy są wywoływane jednocześnie (czasami 2-3 razy dłużej) z powodu jakiegoś problemu z rywalizacją. Nie rozumiem, dlaczego tak musi być, ponieważ przestrzeń alokacji pamięci dla 2 procesorów graficznych jest całkowicie niezależna, więc nie powinno istnieć blokada na cały system cudaMalloc - blokada per-GPU byłaby bardziej uzasadniona.

Aby udowodnić moją hipotezę, że problem dotyczy jednoczesnych wywołań cudaMalloc, stworzyłem niewiarygodnie prosty program z dwoma wątkami CPU (dla każdego GPU), z których każdy wywołuje cudaMalloc kilka razy. Po raz pierwszy prowadził ten program tak, że oddzielne wątki nie nazywają cudaMalloc w tym samym czasie:

Widać to trwa ~ 175 mikrosekund na przydział. Następnie wpadłem program z nitki nazywając cudaMalloc jednocześnie:

Tutaj każde wywołanie trwało ~ 538 mikrosekund lub 3 razy dłuższe niż w poprzednim przypadku! Nie trzeba dodawać, że to bardzo spowalnia moją aplikację i jest oczywiste, że problem będzie tylko gorszy, gdy pojawi się więcej niż 2 GPU.

Zauważyłem to zachowanie w systemach Linux i Windows. W Linuksie używam sterownika Nvidia w wersji 319.60, a w Windows używam wersji 327.23. Używam zestawu narzędzi CUDA 5.5.

Możliwe przyczyny: Używam GTX 690 w tych testach. Ta karta to w zasadzie 2 680-podobne procesory graficzne umieszczone w tej samej jednostce. Jest to jedyna instalacja "multi-GPU", więc może problem z cudaMalloc ma związek z zależnościami sprzętowymi pomiędzy procesorami graficznymi 690?

Answer 1

Podsumowując problemu i otrzymano możliwe rozwiązania

cudaMalloc rywalizacji prawdopodobnie wynika z rywalizacji poziomie sterownika (prawdopodobnie z powodu konieczności przełączania kontekstu urządzenia jak talonmies suggestsed) i można uniknąć dodatkowego opóźnienia w Wykonaj krytyczne sekcje przez cudaMalloc-ing i tymczasowe bufory wcześniej.

Wygląda na to, że prawdopodobnie muszę zmienić kod, aby nie wywoływać żadnej procedury sortowania, która wywołuje cudaMalloc pod maską (w moim przypadku ciąg :: sort_by_key). The CUB library wygląda obiecująco pod tym względem. Jako bonus, CUB udostępnia również użytkownikowi parametr strumienia CUDA, co może również poprawić wydajność.

Aby uzyskać szczegółowe informacje na temat przejścia z ciągu do CUB, patrz CUB (CUDA UnBound) equivalent of thrust::gather.

UPDATE:

I wycofał się z połączenia do oporowej :: sort_by_key na rzecz ow :: DeviceRadixSort :: SortPairs.
Wykonanie tego ogolonego milisekundy poza moim czasem przetwarzania w przedziale czasowym. Rozwiązany został również problem rywalizacji wielu GPU - odciążenie do 2 procesorów graficznych niemal skraca czas przetwarzania o 50%, zgodnie z oczekiwaniami.

Answer 2

Przedmówię to z zastrzeżeniem: nie jestem wtajemniczony do elementów wewnętrznych sterownika NVIDIA, więc jest to nieco spekulatywne.

Spowolnienie, jakie widzisz, to tylko rywalizacja na poziomie kierowców spowodowana konkurencją z wielu wątków, która jednocześnie wywołuje funkcję malloc urządzenia. Przydzielanie pamięci urządzenia wymaga wielu wywołań systemu operacyjnego, podobnie jak przełączanie kontekstu na poziomie sterownika. W obu operacjach występuje niezwyczajna ilość opóźnień. Prawdopodobnie dodatkowy czas, w którym dwa wątki próbują przydzielić pamięć jednocześnie, jest spowodowany dodatkowym opóźnieniem sterownika z przełączania z jednego urządzenia na inne w ramach sekwencji wywołań systemowych wymaganych do przydzielenia pamięci na oba urządzenia.

mogę myśleć kilka sposobów powinieneś być w stanie złagodzić to:

• Można zmniejszyć obciążenie wywołań systemowych naporu alokacji pamięci zeru pisząc swój własny utopił przydzielania pamięci dla urządzenie, które działa z płyty pamięci przydzielonej podczas inicjowania. To pozbyłoby się całego kosztu wywołania systemowego w obrębie każdego sort_by_key, ale wysiłek napisania własnego menedżera pamięci użytkownika jest nie banalny. Z drugiej strony pozostawia resztę kodu ciągłego w nienaruszonym stanie.
• Możesz przełączyć się do alternatywnej biblioteki sortowania i odebrać samodzielnie zarządzać przydzielaniem pamięci tymczasowej. Jeśli wykonasz wszystkie alokacje w fazie inicjalizacji, koszt jednorazowej alokacji pamięci może być zamortyzowany prawie do zera w ciągu każdego wątku.

W wielu algebra liniowych algebry liniowej CUBLAS, które napisałem, połączyłem oba pomysły i napisałem niezależny menedżer pamięci urządzenia użytkownika, który działa z puli przydzielonej pamięci urządzenia raz. Odkryłem, że usunięcie całego kosztu pośredniego alokacji pamięci urządzeń przyniosło użyteczne przyspieszenie. Twój przypadek użycia może skorzystać z podobnej strategii w przypadku.