Rozdział 11: Kierunki badań nad GPU w zakresie skalowania i wykonywania afinicznego
Jak opisano w rozdziale 2, interfejsy programowania GPU, takie jak CUDA i OpenCL, posiadają model programowania podobny do MIMD, który pozwala programiście uruchomić dużą tablicę wątków skalarnych na GPU. Podczas gdy każdy z tych wątków skalarnych może podążać własną unikalną ścieżką wykonania i może uzyskiwać dostęp do dowolnych lokalizacji pamięci, w typowym przypadku wszystkie one podążają za małym zestawem ścieżek wykonania i wykonują podobne operacje.
Zbieżny przepływ sterowania wśród wątków GPU jest wykorzystywany w większości, jeśli nie we wszystkich, nowoczesnych GPU za pośrednictwem modelu wykonania SIMT, w którym wątki skalarne są grupowane w warpy, które działają na sprzęcie SIMD (patrz sekcja 3.1.1). Ten rozdział podsumowuje serię badań, które dalej wykorzystują podobieństwo tych wątków skalarnych za pomocą skalowania i wykonywania afinicznego.
Kluczowe spostrzeżenie tych badań polega na obserwacji struktury wartości [Kim i in., 2013] wśród wątków wykonujących ten sam kernel obliczeniowy. Dwa typy struktury wartości, jednolita i afiniczna, są zilustrowane w kernelu obliczeniowym w przykładzie 11.1.
Zmienna jednolita
Zmienna, która ma tę samą stałą wartość dla każdego wątku w kernelu obliczeniowym. W algorytmie 11.1 zmienna a, a także literały THRESHOLD i Y_MAX_VALUE, mają jednolite wartości we wszystkich wątkach w kernelu obliczeniowym. Zmienna jednolita może być przechowywana w pojedynczym rejestrze skalarnym i używana przez wszystkie wątki w kernelu obliczeniowym.
Zmienna afiniczna
Zmienna, której wartości są liniową funkcją identyfikatora wątku dla każdego wątku w kernelu obliczeniowym. W algorytmie 11.1 adres pamięci zmiennej y[idx] może być reprezentowany jako transformacja afiniczna identyfikatora wątku threadIdx.x:
&(y[idx]) = &(y[0]) + sizeof(int) * threadIdx.x;Ta reprezentacja afiniczna może być przechowywana jako para wartości skalarnych, podstawa i krok, co jest znacznie bardziej zwięzłe niż w pełni rozwinięty wektor.
__global__ void vsadd( int y[], int a ) {
// Kod kernela
}
```Oto polski przekład pliku Markdown:
```c
int idx = threadIdx.x;
y[idx] = y[idx] + a;
if ( y[idx] > PRÓG )
y[idx] = MAKS_WARTOŚĆ_Y;
}Algorytm 11.1: Przykład operacji skalarnych i afinicznych w jądrze obliczeniowym (z [Kim i in., 2013]).
Istnieje wiele propozycji badawczych dotyczących wykrywania i wykorzystywania zmiennych jednolitych lub afinicznych w GPU. Reszta tego rozdziału podsumowuje te propozycje w tych dwóch aspektach.
Wykrywanie zmiennych jednolitych lub afinicznych
Istnieją dwa główne podejścia do wykrywania istnienia zmiennych jednolitych lub afinicznych w jądrze obliczeniowym GPU: Wykrywanie sterowane przez kompilator i Wykrywanie za pomocą sprzętu.
Wykrywanie sterowane przez kompilator
Jednym ze sposobów wykrywania istnienia zmiennych jednolitych lub afinicznych w jądrze obliczeniowym GPU jest przeprowadzenie specjalnej analizy kompilatora. Jest to możliwe, ponieważ istniejące modele programowania GPU, CUDA i OpenCL, już zapewniają środki, aby programista mógł zadeklarować zmienną jako stałą w całym jądrze obliczeniowym, a także zapewniają specjalną zmienną dla identyfikatora wątku. Kompilator może przeprowadzić analizę zależności sterowania, aby wykryć zmienne, które zależą wyłącznie od stałych i identyfikatorów wątków, i oznaczyć je jako jednolite/afiniczne. Operacje działające wyłącznie na zmiennych jednolitych/afinicznych są następnie kandydatami do skalaryzacji.
AMD GCN [AMD, 2012] opiera się na kompilatorze, aby wykryć zmienne jednolite i operacje skalarne, które mogą być przechowywane i przetwarzane przez dedykowany procesor skalarny.
Asanovic i in. [2013] wprowadzają połączoną analizę zbieżną i wariantową, która pozwala kompilatorowi określić operacje w dowolnym jądrze obliczeniowym, które kwalifikują się do skalaryzacji i/lub transformacji afinicznej. Instrukcje w zbieżnych regionach jądra obliczeniowego mogą być przekonwertowane na instrukcje skalarne/afiniczne. W każdym przejściu z regionów rozbieżnych do zbieżnych jądra obliczeniowego kompilator wstawia instrukcję syncwarp, aby obsłużyć zależności rejestrów wywołane przez przepływ sterowania między tymi regionami. Asanovic i in. [2013] przyjęli tOto tłumaczenie na język polski:
Niniejsza analiza ma na celu wygenerowanie operacji skalarnych dla architektury Temporal-SIMT [Keckler i in., 2011, Krashinsky, 2011].
Oddzielone Obliczenia Afiniczne (DAC) [Wang i Lin, 2017] opierają się na podobnej analizie kompilatora, aby wyodrębnić kandydatów skalarnych i afinicznych, którzy mają być oddzieleni do osobnego warpu. Wang i Lin [2017] wzbogacają ten proces o analizę rozbieżnych afinicznych, mając na celu wyodrębnienie pasm instrukcji