Jak są zorganizowane wątki, które mają być wykonywane przez GPU?
Jak są zorganizowane wątki, które mają być wykonywane przez GPU?
Odpowiedzi:
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).
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).
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;
Załóżmy, że GPU 9800GT:
https://www.tutorialspoint.com/cuda/cuda_threads.htm
Blok nie może mieć więcej aktywnych wątków niż 512, dlatego __syncthreads
moż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:
Uwaga:
Głównym punktem jest __syncthreads
operacja obejmująca cały blok i nie synchronizuje ona wszystkich wątków.
Nie jestem pewien dokładnej liczby wątków, które __syncthreads
moż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.
__syncthreads
jest 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.