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

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

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;

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.