Zrozumienie wymiarów siatki CUDA, wymiarów bloków i organizacji wątków (proste wyjaśnienie) [zamknięte]

161

Jak są zorganizowane wątki, które mają być wykonywane przez GPU?

cibercitizen1
źródło
Przewodnik programowania CUDA powinien być dobrym miejscem do rozpoczęcia tego. Poleciłbym również zapoznać się z wprowadzeniem do CUDA stąd .
Tom

Odpowiedzi:

287

Sprzęt komputerowy

Jeśli urządzenie GPU ma np. 4 jednostki wieloprocesorowe, a każda z nich może obsługiwać 768 wątków: to w danym momencie nie więcej niż 4 * 768 wątków będzie naprawdę działać równolegle (jeśli zaplanowałeś więcej wątków, będą czekać ich ruch).

Oprogramowanie

wątki są zorganizowane w bloki. Blok jest wykonywany przez jednostkę wieloprocesorową. Wątki bloku mogą być identyfikowane (indeksowane) za pomocą indeksów 1Dimension (x), 2Dimensions (x, y) lub 3Dim (x, y, z), ale w każdym przypadku dla naszego przykładu x y z <= 768 (obowiązują inne ograniczenia do x, y, z, patrz przewodnik i opis możliwości urządzenia).

Oczywiście, jeśli potrzebujesz więcej niż tych 4 * 768 wątków, potrzebujesz więcej niż 4 bloki. Bloki mogą być również indeksowane 1D, 2D lub 3D. Na wejście do GPU czeka kolejka bloków (ponieważ w naszym przykładzie GPU ma 4 procesory i tylko 4 bloki są wykonywane jednocześnie).

Teraz prosty przypadek: przetwarzanie obrazu 512x512

Załóżmy, że chcemy, aby jeden wątek przetwarzał jeden piksel (i, j).

Możemy użyć bloków po 64 wątki każdy. Następnie potrzebujemy 512 * 512/64 = 4096 bloków (aby mieć 512x512 wątków = 4096 * 64)

Często organizuje się (aby ułatwić indeksowanie obrazu) wątki w blokach 2D, które mają blockDim = 8 x 8 (64 wątki na blok). Wolę nazywać to ThreadPerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

i 2D gridDim = 64 x 64 bloki (potrzebne 4096 bloków). Wolę nazywać to numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y); 

Jądro jest uruchamiane w następujący sposób:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );       

Wreszcie: pojawi się coś w rodzaju „kolejki 4096 bloków”, w której blok czeka na przypisanie jednego z procesorów GPU do wykonania 64 wątków.

W jądrze piksel (i, j), który ma zostać przetworzony przez wątek, jest obliczany w następujący sposób:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;
cibercitizen1
źródło
11
Jeśli każdy blok może obsługiwać 768 wątków, po co używać tylko 64? Jeśli użyjesz maksymalnego limitu 768, będziesz mieć mniej bloków, a więc lepszą wydajność.
Aliza
10
@Aliza: bloki są logiczne , limit 768 wątków przypada na każdą fizyczną jednostkę przetwarzającą. Używasz bloków zgodnie ze specyfikacją swojego problemu, aby rozdzielić pracę na wątki. Nie jest prawdopodobne, aby zawsze można było użyć bloków 768 wątków dla każdego problemu. Wyobraź sobie, że musisz przetworzyć obraz 64x64 (4096 pikseli). 4096/768 = 5,333333 bloki?
cibercitizen1
1
blok są logiczne, ale każdy blok jest przypisany do rdzenia. jeśli bloków jest więcej niż rdzeni, bloki są ustawiane w kolejce do zwolnienia rdzeni. W twoim przykładzie możesz użyć 6 bloków, a dodatkowe wątki nic nie robią (2/3 wątków w szóstym bloku).
Aliza
3
@ cibercitizen1 - Myślę, że uwaga Alizy jest dobra: jeśli to możliwe, chce się użyć jak największej liczby wątków na blok. Jeśli istnieje ograniczenie, które wymaga mniejszej liczby wątków, lepiej wyjaśnij, dlaczego może tak być w drugim przykładzie (ale najpierw wyjaśnij prostszy i bardziej pożądany przypadek).
6
@thouis Tak, może. Ale przypadek jest taki, że ilość pamięci potrzebnej dla każdego wątku zależy od aplikacji. Na przykład w moim ostatnim programie każdy wątek wywołuje funkcję optymalizacji metodą najmniejszych kwadratów, wymagającą „dużo” pamięci. Do tego stopnia, że ​​bloki nie mogą być większe niż nici 4x4. Mimo to uzyskane przyspieszenie było dramatyczne w porównaniu z wersją sekwencyjną.
cibercitizen1
9

Załóżmy, że GPU 9800GT:

  • posiada 14 multiprocesorów (SM)
  • każdy SM ma 8 procesorów wątkowych (procesory strumieniowe AKA, SP lub rdzenie)
  • umożliwia do 512 wątków na blok
  • warpsize wynosi 32 (co oznacza, że ​​każdy z procesorów wątków 14x8 = 112 może zaplanować do 32 wątków)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Blok nie może mieć więcej aktywnych wątków niż 512, dlatego __syncthreadsmoże synchronizować tylko ograniczoną liczbę wątków. tj. jeśli wykonasz następujące czynności z 600 wątkami:

func1();
__syncthreads();
func2();
__syncthreads();

wtedy jądro musi działać dwukrotnie, a kolejność wykonywania będzie następująca:

  1. func1 jest wykonywana dla pierwszych 512 wątków
  2. func2 jest wykonywana dla pierwszych 512 wątków
  3. func1 jest wykonywana dla pozostałych wątków
  4. func2 jest wykonywana dla pozostałych wątków

Uwaga:

Głównym punktem jest __syncthreadsoperacja obejmująca cały blok i nie synchronizuje ona wszystkich wątków.


Nie jestem pewien dokładnej liczby wątków, które __syncthreadsmożna zsynchronizować, ponieważ możesz utworzyć blok z więcej niż 512 wątkami i pozwolić warpowi zająć się planowaniem. W moim rozumieniu dokładniejsze jest stwierdzenie: func1 jest wykonywana co najmniej dla pierwszych 512 wątków.

Zanim zredagowałem tę odpowiedź (w 2010 roku) zmierzyłem, że wątki 14x8x32 zostały zsynchronizowane za pomocą __syncthreads.

Byłbym bardzo wdzięczny, gdyby ktoś przetestował to ponownie w celu uzyskania dokładniejszych informacji.

Bizhan
źródło
Co się stanie, jeśli funkcja func2 () zależy od wyników funkcji func1 (). Myślę, że to źle
Chris
@Chris Napisałem to siedem lat temu, ale jeśli dobrze pamiętam, zrobiłem test na tym i doszedłem do wniosku, że jądra z większą liczbą wątków niż gpu zachowują się w ten sposób. Jeśli zdarzy ci się przetestować ten przypadek i osiągniesz inny wynik, będę musiał usunąć ten post.
Bizhan
Przepraszam, myślę, że to źle, także, że GPU może jednocześnie uruchamiać tylko 112 wątków.
Steven Lu
@StevenLu próbowałeś tego? również nie sądzę, aby 112 współbieżnych wątków miało sens dla GPU. 112 to liczba procesorów strumieniowych. Prawie nie pamiętam CUDA :)
Bizhan
1
@StevenLu maksymalna liczba wątków nie jest tutaj problemem, __syncthreadsjest operacją obejmującą cały blok, a fakt, że w rzeczywistości nie synchronizuje wszystkich wątków, jest uciążliwy dla osób uczących się CUDA. Więc zaktualizowałem moją odpowiedź na podstawie informacji, które mi podałeś. Bardzo to doceniam.
Bizhan