Jakie problemy nadają się dobrze do obliczeń na GPU?

84

Mam więc przyzwoitą uwagę, z jakimi problemami, z którymi pracuję, najlepiej rozwiązać szeregowo, i którymi można zarządzać równolegle. Ale w tej chwili nie mam pojęcia, co najlepiej poradzić sobie z obliczeniami opartymi na procesorze i co powinno zostać odciążone na GPU.

Wiem, że jest to podstawowe pytanie, ale większość moich poszukiwań zostaje uwięzionych w ludziach, którzy wyraźnie opowiadają się za jednym lub drugim, bez uzasadnienia dlaczego , lub nieco niejasnych zasad. Szukasz bardziej użytecznej odpowiedzi tutaj.

Fomite
źródło

Odpowiedzi:

63

Sprzęt GPU ma dwie szczególne zalety: surowe obliczenia (FLOP) i przepustowość pamięci. Najtrudniejsze problemy obliczeniowe należą do jednej z tych dwóch kategorii. Na przykład gęsta algebra liniowa (A * B = C lub Rozwiąż [Ax = y] lub Przekątna [A] itp.) Przypada gdzieś na widmo przepustowości obliczeniowej / pamięci w zależności od wielkości systemu. Szybkie transformaty Fouriera (FFT) również pasują do tej formy o wysokich potrzebach w zakresie łącznej przepustowości. Podobnie jak inne transformacje, algorytmy oparte na siatce / siatce, Monte Carlo itp. Jeśli spojrzysz na przykłady kodu NVIDIA SDK , możesz poczuć, jakie problemy są najczęściej rozwiązywane.

Wydaje mi się, że bardziej pouczającą odpowiedzią jest pytanie: „Na jakich problemach układy GPU są naprawdę złe?” Większość problemów, które nie należą do tej kategorii, można uruchomić na GPU, choć niektóre wymagają więcej wysiłku niż inne.

Problemy, które nie są dobrze odwzorowane, są zazwyczaj zbyt małe lub zbyt nieprzewidywalne. Bardzo małym problemom brakuje równoległości potrzebnej do wykorzystania wszystkich wątków na GPU i / lub mogłyby zmieścić się w pamięci podręcznej niskiego poziomu procesora, znacznie zwiększając wydajność procesora. Nieprzewidywalne problemy mają zbyt wiele znaczących gałęzi, które mogą uniemożliwić wydajne przesyłanie danych z pamięci GPU do rdzeni lub zmniejszyć równoległość poprzez złamanie paradygmatu SIMD (patrz „ rozbieżne wypaczenia ”). Przykłady tego rodzaju problemów obejmują:

  • Większość algorytmów graficznych (zbyt nieprzewidywalnych, szczególnie w przestrzeni pamięci)
  • Rzadka algebra liniowa (ale to również źle na CPU)
  • Małe problemy z przetwarzaniem sygnału (na przykład FFT mniejsze niż 1000 punktów)
  • Szukaj
  • Sortować
Max Hutchinson
źródło
3
Mimo to rozwiązania GPU dla tych „nieprzewidywalnych” problemów możliwe i chociaż obecnie nie są zazwyczaj wykonalne, mogą zyskać na znaczeniu w przyszłości.
lewo około
6
Chciałbym specjalnie dodać gałęzie do listy wyłączników wydajności GPU. Chcesz, aby wszystkie (setki) wykonywały tę samą instrukcję (jak w SIMD), aby wykonać naprawdę równoległe obliczenia. Na przykład na kartach AMD, jeśli którykolwiek z przepływów instrukcji napotyka gałąź i musi się rozchodzić - wszystkie fale frontowe (grupa równoległa) są rozbieżne. Jeśli inne jednostki z frontu fali nie mogą się rozejść - muszą wykonać drugie przejście. Chyba właśnie to oznacza maxhutch przez przewidywalność.
Violet Giraffe
2
@ VioletGiraffe, to niekoniecznie prawda. W CUDA (tj. Na procesorach graficznych Nvidia) rozbieżność gałęzi wpływa tylko na bieżące wypaczenie, które ma maksymalnie 32 wątki. Różne wypaczenia, chociaż wykonują ten sam kod, nie są synchroniczne, chyba że są wyraźnie zsynchronizowane (np. Z __synchtreads()).
Pedro
1
@Pedro: To prawda, ale rozgałęzienie ogólnie szkodzi wydajności. W przypadku kodów o wysokiej wydajności (czym nie jest kod GPU?) Należy wziąć to pod uwagę.
jvriesem,
21

Problemy, które mają wysoką intensywność arytmetyczną i regularne wzorce dostępu do pamięci, są zazwyczaj łatwe do wykonania na GPU i dobrze na nich działają.

Podstawową trudnością w posiadaniu wysokowydajnego kodu GPU jest to, że masz mnóstwo rdzeni i chcesz, aby wszystkie były maksymalnie wykorzystywane. Problemy, które mają nieregularne wzorce dostępu do pamięci lub nie mają dużej intensywności arytmetycznej, utrudniają to: albo spędzasz dużo czasu na komunikowaniu wyników, albo spędzasz dużo czasu na pobieraniu rzeczy z pamięci (co jest powolne!), I za mało czasu na zgniatanie liczb. Oczywiście potencjał współbieżności w twoim kodzie jest krytyczny dla jego zdolności do dobrego zaimplementowania również na GPU.

Reid.Atcheson
źródło
Czy możesz określić, co masz na myśli przez regularne wzorce dostępu do pamięci?
Fomite
1
Odpowiedź maxhutch jest lepsza niż moja. Rozumiem przez regularny wzorzec dostępu, że dostęp do pamięci jest czasowy i przestrzenny lokalnie. To znaczy: nie wykonujesz wielkich skoków wokół pamięci wielokrotnie. To także coś w rodzaju pakietu, który zauważyłem. Rozumie się przez to, że wzorce dostępu do danych mogą być z góry określone albo przez kompilator, albo przez programistę, aby zminimalizować rozgałęzienia (instrukcje warunkowe w kodzie).
Reid.Atcheson
15

Nie jest to samodzielna odpowiedź, ale dodatek do innych odpowiedzi autorstwa maxhutch i Reid.Atcheson .

Aby maksymalnie wykorzystać możliwości procesorów graficznych, Twój problem musi być nie tylko wysoce (lub masowo) równoległy, ale także podstawowy algorytm, który będzie wykonywany na GPU, powinien być jak najmniejszy. W terminologii OpenCL jest to najczęściej nazywane jądrem .

Mówiąc ściślej, jądro powinno pasować do rejestru każdej jednostki wieloprocesowej (lub jednostki obliczeniowej ) GPU. Dokładny rozmiar rejestru zależy od procesora graficznego.

Ze względu na jądro jest na tyle mały, surowe dane problemu musi pasować do pamięci lokalnej GPU (czytaj: pamięć lokalną (OpenCL) lub pamięci współdzielonej (CUDA) jednostki obliczeniowej). W przeciwnym razie nawet wysoka przepustowość pamięci GPU nie jest wystarczająco szybka, aby cały czas przetwarzać elementy przetwarzające .
Zazwyczaj ta pamięć jest około 16 do 32 KiByte duży .

Torbjörn
źródło
Czy pamięć lokalna / współużytkowana każdej jednostki przetwarzania nie jest współużytkowana przez wszystkie dziesiątki (?) Wątków działających w jednym klastrze rdzeni? W takim przypadku, czy naprawdę nie musisz znacznie zmniejszać roboczego zestawu danych, aby uzyskać pełną wydajność GPU?
Dan Neely
Lokalna / współdzielona pamięć jednostki przetwarzającej jest dostępna tylko dla samej jednostki obliczeniowej, a zatem współużytkowana tylko przez elementy przetwarzające tej jednostki obliczeniowej. Globalna pamięć karty graficznej (zwykle 1 GB) jest dostępna dla wszystkich jednostek przetwarzających. Przepustowość między elementami przetwarzającymi a pamięcią lokalną / współużytkowaną jest bardzo szybka (> 1 TB / s), ale przepustowość do pamięci globalnej jest znacznie wolniejsza (~ 100 GB / s) i musi być dzielona między wszystkie jednostki obliczeniowe.
Torbjörn
Nie pytałem o główną pamięć GPU. Myślałem, że pamięć na matrycy została przydzielona tylko na poziomie rdzenia, a nie na każdy rdzeń. np. dla nVidia GF100 / 110 gpu; dla każdego z 16 klastrów SM, a nie 512 rdzeni cuda. Z każdym SM zaprojektowanym do równoległego działania do 32 wątków maksymalizacja wydajności GPU wymagałaby utrzymania zestawu roboczego w zakresie 1 kb / wątku.
Dan Neely
@Torbjoern Chcemy, aby wszystkie procesory wykonawcze GPU były zajęte, procesory graficzne osiągają to na dwa sposoby: (1) najczęstszym sposobem jest zwiększenie zajętości, lub inaczej mówiąc, poprzez zwiększenie liczby współbieżnych wątków (małe jądra zużywają mniej współdzielone zasoby, dzięki czemu możesz mieć więcej aktywnych wątków); może lepiej jest (2) zwiększyć równoległość poziomu instrukcji w twoim jądrze, abyś mógł mieć większe jądro przy stosunkowo niskim obłożeniu (mała liczba aktywnych wątków). Zobacz bit.ly/Q3KdI0
fcruz,
11

Prawdopodobnie bardziej techniczny dodatek do poprzednich odpowiedzi: Procesory graficzne CUDA (tj. Nvidia) można opisać jako zestaw procesorów, które działają niezależnie na 32 wątkach. Wątki w każdym procesorze działają w trybie blokowania (pomyśl SIMD z wektorami o długości 32).

Chociaż najbardziej kuszącym sposobem pracy z procesorami graficznymi jest udawanie, że absolutnie wszystko działa w trybie blokowania, nie zawsze jest to najbardziej efektywny sposób robienia rzeczy.

Jeśli twój kod nie ładnie / automatycznie łączy się równolegle z setkami / tysiącami wątków, możesz być w stanie rozbić go na pojedyncze zadania asynchroniczne, które dobrze się zrównoleglają i wykonać te z tylko 32 wątkami działającymi w trybie blokowania. CUDA zapewnia zestaw instrukcji atomowych, które umożliwiają implementację muteksów, co z kolei umożliwia procesorom synchronizację między sobą i przetwarzanie listy zadań w paradygmacie puli wątków . Twój kod działałby wtedy w podobny sposób jak w systemie wielordzeniowym, pamiętaj tylko, że każdy rdzeń ma wtedy 32 własne wątki.

Oto mały przykład zastosowania CUDA, jak to działa

/* Global index of the next available task, assume this has been set to
   zero before spawning the kernel. */
__device__ int next_task;

/* We will use this value as our mutex variable. Assume it has been set to
   zero before spawning the kernel. */
__device__ int tasks_mutex;

/* Mutex routines using atomic compare-and-set. */
__device__ inline void cuda_mutex_lock ( int *m ) {
    while ( atomicCAS( m , 0 , 1 ) != 0 );
    }
__device__ inline void cuda_mutex_unlock ( int *m ) {
    atomicExch( m , 0 );
    }

__device__ void task_do ( struct task *t ) {

    /* Do whatever needs to be done for the task t using the 32 threads of
       a single warp. */
    }

__global__ void main ( struct task *tasks , int nr_tasks ) {

    __shared__ task_id;

    /* Main task loop... */
    while ( next_task < nr_tasks ) {

        /* The first thread in this block is responsible for picking-up a task. */
        if ( threadIdx.x == 0 ) {

            /* Get a hold of the task mutex. */
            cuda_mutex_lock( &tasks_mutex );

            /* Store the next task in the shared task_id variable so that all
               threads in this warp can see it. */
            task_id = next_task;

            /* Increase the task counter. */
            next_tast += 1;

            /* Make sure those last two writes to local and global memory can
               be seen by everybody. */
            __threadfence();

            /* Unlock the task mutex. */
            cuda_mutex_unlock( &tasks_mutex );

            }

        /* As of here, all threads in this warp are back in sync, so if we
           got a valid task, perform it. */
        if ( task_id < nr_tasks )
            task_do( &tasks[ task_id ] );

        } /* main loop. */

    }

Następnie musisz wywołać jądro, main<<<N,32>>>(tasks,nr_tasks)aby upewnić się, że każdy blok zawiera tylko 32 wątki, a tym samym mieści się w jednym warp. W tym przykładzie założyłem również, dla uproszczenia, że ​​zadania nie mają żadnych zależności (np. Jedno zadanie zależy od wyników innego) lub konfliktów (np. Praca na tej samej pamięci globalnej). W takim przypadku wybór zadania staje się nieco bardziej skomplikowany, ale struktura jest zasadniczo taka sama.

Jest to oczywiście bardziej skomplikowane niż robienie wszystkiego na jednej dużej partii komórek, ale znacznie poszerza rodzaj problemów, do których można użyć procesorów graficznych.

Pedro
źródło
2
Jest to technicznie prawdą, ale wymagana jest wysoka równoległość, aby uzyskać wysoką przepustowość pamięci i istnieje ograniczenie liczby asynchronicznych wywołań jądra (obecnie 16). Tobie jest także mnóstwo nieudokumentowanych zachowań związanych z planowaniem w bieżącej wersji. Odradzam poleganie na asynchronicznych jądrach w celu poprawy wydajności na razie ...
Max Hutchinson
2
To, co opisuję, można zrobić w jednym wywołaniu jądra. Możesz utworzyć N bloków po 32 wątki, tak aby każdy blok pasował do jednej warp. Każdy blok następnie pobiera zadanie z globalnej listy zadań (kontrolowanej za pomocą atomics / mutexes) i oblicza je za pomocą 32 wątków z blokadą. Wszystko to dzieje się w jednym wywołaniu jądra. Jeśli chcesz przykład kodu, daj mi znać, a ja go opublikuję.
Pedro
4

Jak dotąd nie stwierdzono, że obecna generacja układów GPU nie radzi sobie tak dobrze w obliczeniach zmiennoprzecinkowych podwójnej precyzji, jak w obliczeniach pojedynczej precyzji. Jeśli obliczenia muszą być wykonywane z podwójną precyzją, można oczekiwać, że czas działania wzrośnie 10-krotnie w stosunku do pojedynczej precyzji.

Brian Borchers
źródło
Chcę się nie zgodzić. Większość (lub wszystkie) nowsze karty graficzne mają natywną obsługę podwójnej precyzji. Prawie każdy taki GPU zgłasza obliczenia obliczeń z podwójną precyzją działające z około połową prędkości pojedynczej precyzji, prawdopodobnie z powodu prostego podwojenia wymaganego dostępu do pamięci / przepustowości.
Godric Seer
1
Chociaż prawdą jest, że najnowsze i najlepsze karty Nvidia Tesla oferują najwyższą wydajność podwójnej precyzji, która stanowi połowę szczytowej wydajności pojedynczej precyzji, stosunek ten wynosi 8 do 1 w przypadku bardziej powszechnych kart konsumenckich z architekturą Fermi.
Brian Borchers,
@GodricSeer Współczynnik zmiennoprzecinkowy SP i DP 2: 1 ma bardzo niewiele wspólnego z przepustowością i prawie wszystko z tym, ile jednostek sprzętowych wykonuje te operacje. Powszechne jest ponowne użycie pliku rejestru dla SP i DP, dlatego jednostka zmiennoprzecinkowa może wykonać 2x operacje SP jako operacje DP. Istnieje wiele wyjątków od tego projektu, np. IBM Blue Gene / Q (nie ma logiki SP, a zatem SP działa przy ~ 1.05x DP). Niektóre GPU mają współczynniki inne niż 2, np. 3 i 5.
Jeff
Minęły cztery lata, odkąd napisałem tę odpowiedź, a obecna sytuacja z procesorami graficznymi NVIDIA jest taka, że ​​w przypadku linii GeForce i Quadro stosunek DP / SP wynosi teraz 1/32. Procesory graficzne NVIDIA Tesla mają znacznie lepszą wydajność podwójnej precyzji, ale także kosztują znacznie więcej. Z drugiej strony AMD nie obniżyło wydajności podwójnej precyzji na swoich kartach graficznych Radeon w ten sam sposób.
Brian Borchers
4

Z metaforycznego punktu widzenia, gpu można postrzegać jako osobę leżącą na łóżku z paznokci. Osoba leżąca na górze to dane, a u podstawy każdego gwoździa znajduje się procesor, więc gwóźdź jest w rzeczywistości strzałką wskazującą od procesora do pamięci. Wszystkie paznokcie mają regularny wzór, jak siatka. Jeśli ciało jest dobrze rozłożone, czuje się dobrze (wydajność jest dobra), jeśli ciało dotyka tylko niektórych miejsc łożyska paznokcia, wtedy ból jest zły (zła wydajność).

Można to uznać za komplementarną odpowiedź na powyższe doskonałe odpowiedzi.

labotsirc
źródło
4

Stare pytanie, ale myślę, że ta odpowiedź z 2014 r. - związana z metodami statystycznymi, ale możliwa do uogólnienia dla każdego, kto wie, czym jest pętla - jest szczególnie ilustracyjna i pouczająca.

GT.
źródło
2

Procesory graficzne mają długi czas oczekiwania we / wy, więc do nasycenia pamięci trzeba użyć wielu wątków. Aby warp był zajęty, wymaga wielu wątków. Jeśli ścieżka kodu wynosi 10 zegarów, a opóźnienie we / wy 320 zegarów, 32 wątki powinny zbliżyć się do nasycenia warp. Jeśli ścieżka do kodu wynosi 5 zegarów, należy podwoić wątki.

Z tysiącem rdzeni szukaj tysięcy wątków, aby w pełni wykorzystać procesor graficzny.

Dostęp do pamięci odbywa się za pomocą linii pamięci podręcznej, zwykle 32 bajty. Ładowanie jednego bajtu ma koszt porównywalny do 32 bajtów. Połącz więc pamięć, aby zwiększyć lokalność użytkowania.

Każda osnowa zawiera wiele rejestrów i lokalnej pamięci RAM, co umożliwia dzielenie się przez sąsiadów.

Symulacje bliskości dużych zestawów powinny dobrze się optymalizować.

Losowe we / wy i pojedyncze wątki to zabójcza radość ...

użytkownik14381
źródło
To jest naprawdę fascynujące pytanie; Kłócę się ze sobą, czy możliwe jest (lub warte wysiłku) „równoległe” racjonalnie proste zadanie (wykrywanie krawędzi na zdjęciach lotniczych), gdy każde zadanie zajmuje ~ 0,06 s, ale do wykonania jest ~ 1,8 miliona zadań ( rocznie, dla danych o wartości 6 lat: zadania są zdecydowanie rozdzielne) ... a zatem ~ 7,5 dnia obliczeniowego czasu na jednym rdzeniu. Jeśli każde obliczenie było szybsze na GPU, a zadanie można sparaliżować 1-na-nGPUcores [n małe], to czy rzeczywiście jest prawdopodobne, że czas zadania może spaść do ~ 1 godziny? Wydaje się mało prawdopodobne.
GT.
0

Wyobraź sobie problem, który można rozwiązać za pomocą brutalnej siły, jak na przykład Traveling Salesman. Wyobraź sobie, że masz szafy serwerów z 8 klarownymi kartami wideo, a każda karta ma 3000 rdzeni CUDA.

Wystarczy rozwiązać WSZYSTKIE możliwe trasy sprzedawcy, a następnie posortować według czasu / odległości / niektórych danych. Jasne, że wyrzucasz prawie 100% swojej pracy, ale brutalna siła jest czasem realnym rozwiązaniem.

Criggie
źródło
Miałem dostęp do małej farmy 4 takich serwerów przez tydzień, aw ciągu pięciu dni zrobiłem więcej bloków Distributed.net niż w ciągu ostatnich 10 lat.
Criggie
-1

Po przestudiowaniu wielu pomysłów inżynieryjnych powiedziałbym, że GPU jest formą skupiania się zadań, zarządzania pamięcią, powtarzalnych obliczeń.

Wiele formuł może być prostych do napisania, ale bolesnych do obliczenia, na przykład w matematyce matematycznej nie otrzymujesz jednej odpowiedzi, ale wiele wartości.

Jest to ważne w obliczeniach, ponieważ komputer oblicza wartości i uruchamia formuły, ponieważ niektóre formuły nie mogą działać bez wszystkich obliczonych wartości (dlatego zwalniają). Komputer nie bardzo dobrze wie, w jakiej kolejności uruchamiać formuły lub obliczać wartości do użycia w tych programach. Głównie przebija siły przy dużych prędkościach i rozbija formuły na uchwyty, aby je obliczyć, ale wiele programów w dzisiejszych czasach wymaga tych obliczonych uchwytów teraz i czeka na pytania (i pytania i więcej pytań).

Na przykład w grze symulacyjnej, która powinna być obliczana najpierw w zderzeniach, uszkodzenie kolizji, położenie obiektów, nowa prędkość? Ile czasu to zajmie? Jak każdy procesor może poradzić sobie z tym obciążeniem? Ponadto większość programów jest bardzo abstrakcyjna i wymaga więcej czasu na przetwarzanie danych i nie zawsze jest zaprojektowana do wielowątkowości lub nie ma dobrych sposobów na efektywne wykonywanie takich programów w programach abstrakcyjnych.

Gdy procesor stał się lepszy, a lepsi ludzie stali się niechlujni w programowaniu i musimy programować również dla wielu różnych typów komputerów. GPU zaprojektowano tak, aby brutalnie wykorzystywało wiele prostych obliczeń w tym samym czasie (nie wspominając o pamięci (wtórnej / pamięci RAM), a chłodzenie przez ogrzewanie to główne szyjki butelek w komputerach). Jednostka centralna zarządza wieloma pytaniami jednocześnie lub jest wciągana w wiele kierunków, zastanawia się, czego nie można zrobić. (hej, to prawie ludzkie)

GPU jest cholernym pracownikiem nużącym dziełem. Procesor zarządza całkowitym chaosem i nie jest w stanie poradzić sobie z każdym szczegółem.

Czego się uczymy? GPU wykonuje szczegółową żmudną pracę naraz, a procesor jest maszyną wielozadaniową, która nie potrafi się dobrze skoncentrować przy zbyt dużej liczbie zadań do wykonania. (To tak, jakby miało jednocześnie zaburzenie uwagi i autyzm).

Inżynieria to pomysły, design, rzeczywistość i dużo cholernej roboty.

Kiedy wychodzę, pamiętaj, aby zacząć od razu, zacznij szybko, szybko, szybko i szybko i nigdy nie przestawaj próbować.

Andrew G. Corbi
źródło