Rozdział 7: Projektowanie Wieloprocesorów Strumieniowych w Architekturze GPU

Wieloprocesor strumieniowy (SM) jest podstawowym blokiem budulcowym architektur GPU firmy NVIDIA. Każdy SM zawiera zestaw rdzeni CUDA, które wykonują instrukcje w stylu SIMT (Single Instruction, Multiple Thread). SM odpowiada za zarządzanie i planowanie warptów, obsługę rozbieżności gałęzi oraz zapewnienie szybkiego dostępu do pamięci współdzielonej i pamięci podręcznych. W tym rozdziale zbadamy mikroarchitekturę SM, w tym jej potoki, mechanizmy planowania warptów, projekt pliku rejestru oraz organizację pamięci współdzielonej i pamięci podręcznej L1.

Mikroarchitektura i Potoki SM

SM jest wysoce równoległym i potokowanym procesorem zaprojektowanym do wydajnego wykonywania setek wątków równolegle. Rysunek 7.1 przedstawia uproszczony schemat blokowy SM w architekturze NVIDIA Volta.

                                 Pamięć podręczna instrukcji
                                         |
                                         v
                                    Planista Warptów
                                         |
                                         v
                               Jednostka Dyspozycji (4 warpty)
                                 |   |   |   |
                                 v   v   v   v
                               Rdzeń CUDA (FP64/FP32/INT)
                               Rdzeń CUDA (FP64/FP32/INT)
                               Rdzeń CUDA (FP64/FP32/INT)
                               ...
                               Rdzeń Tensorowy
                               Rdzeń Tensorowy
                               ...
                               Jednostka Ładowania/Przechowywania
                               Jednostka Ładowania/Przechowywania
                               ...
                               Jednostka Funkcji Specjalnych
                                         ^
                                         |
                                Plik Rejestrów (64 KB)
                                         ^
```Tutaj jest polska wersja tłumaczenia pliku Markdown z komentarzami przetłumaczonymi na język polski, a kodem pozostawionym bez zmian:

                                  Pamięć współdzielona / Pamięć podręczna L1 (96 KB)

Rysunek 7.1: Uproszczony diagram blokowy SM w architekturze NVIDIA Volta.

Główne komponenty SM obejmują:

Pamięć podręczna instrukcji: Przechowuje często używane instrukcje w celu zmniejszenia opóźnienia i poprawy wydajności.
Planista warp: Wybiera warpy gotowe do wykonania i wysyła je do dostępnych jednostek wykonawczych.
Jednostka dyspozycji: Pobiera i dekoduje instrukcje dla maksymalnie 4 warpów na cykl i wysyła je do odpowiednich jednostek wykonawczych.
Rdzenie CUDA: Programowalne jednostki wykonawcze obsługujące szeroką gamę operacji całkowitoliczbowych i zmiennoprzecinkowych. Każdy SM w Volta zawiera 64 rdzenie CUDA.
Rdzenie Tensor: Wyspecjalizowane jednostki wykonawcze zaprojektowane do przyspieszania obciążeń uczenia głębokiego i AI. Każdy SM w Volta zawiera 8 rdzeni Tensor.
Jednostki ładowania/przechowywania: Obsługują operacje pamięciowe, w tym ładowanie i przechowywanie w pamięci globalnej, pamięci współdzielonej i pamięciach podręcznych.
Jednostki funkcji specjalnych: Wykonują działania trygonometryczne i inne złożone operacje matematyczne.
Rejestr plików: Zapewnia szybki dostęp do rejestrów prywatnych wątków. Każdy SM w Volta ma 64 KB rejestru plików.
Pamięć współdzielona / Pamięć podręczna L1: Konfigurowalny obszar pamięci, który może być używany jako pamięć podręczna zarządzana przez oprogramowanie (pamięć współdzielona) lub jako sprzętowo zarządzana pamięć podręczna danych L1.

Potok SM jest zaprojektowany w celu maksymalizacji przepustowości, umożliwiając jednoczesne wykonywanie wielu warpów i ukrywanie opóźnień pamięci. Rysunek 7.2 ilustruje uproszczony widok potoku SM.

    Pobranie instrukcji
            |
            v
    Dekodowanie instrukcji
            |
            v
    Zbieranie operandów
            |
            v
    Wykonywanie (rdzenie CUDA, rdzenie Tensor, jednostki ładowania/przechowywania, jednostki funkcji specjalnych)
            |
            v
    Zapis wyników

Rysunek 7.2: Uproszczony potok SM.

Etapy potoku to:

Pobranie instrukcji: Planista warpów wybiera warp gotowy do wykonania i wysyła go do jednostki dyspozycji.


1. **Pobranie instrukcji**: Procesor SM pobiera następną instrukcję dla tej warpy z pamięci cache instrukcji.

2. **Dekodowanie instrukcji**: Pobrana instrukcja jest dekodowana, aby określić typ operacji, operandy i rejestry docelowe.

3. **Zebranie operandów**: Wymagane operandy dla tej instrukcji są pobierane z rejestru lub pamięci współdzielonej.

4. **Wykonanie**: Instrukcja jest wykonywana na odpowiedniej jednostce wykonawczej (CUDA Core, Tensor Core, Jednostka ładowania/przechowywania lub Jednostka Funkcji Specjalnych).

5. **Zapis wyniku**: Wynik wykonania jest zapisywany z powrotem do rejestru lub pamięci współdzielonej.

Aby osiągnąć wysoką wydajność, SM wykorzystuje kilka technik w celu zmaksymalizowania wykorzystania zasobów i ukrycia opóźnień:

- **Podwójne wydawanie**: SM może wydawać dwie niezależne instrukcje na warpę w pojedynczym cyklu, co pozwala na zwiększenie równoległości na poziomie instrukcji.
- **Spipelinowane jednostki wykonawcze**: Jednostki wykonawcze są spipelinowane, co umożliwia SM rozpoczęcie nowej operacji na jednostce przed zakończeniem poprzedniej.
- **Ukrywanie opóźnień**: SM może przełączać się między warpami w cyklu taktowania, pozwalając na ukrycie opóźnień dostępu do pamięci i operacji o długim czasie oczekiwania poprzez wykonywanie instrukcji z innych warp.

Przykład 7.1 pokazuje prosty kernel CUDA, który wykonuje elementarną operację dodawania dwóch wektorów.

```cpp
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Przykład 7.1: Kernel CUDA dla dodawania wektorów.

W tym przykładzie, każdy wątek w kernelu oblicza sumę odpowiadających sobie elementów z wektorów wejściowych a i b i zapisuje wynik w wektorze wyjściowym c. SM wykonuje ten kernel poprzez przypisanie każdego wątku do rdzenia CUDA i planowanie warp wątków do wykonania na dostępnych rdzeniach. Jednostki ładowania/przechowywania są wykorzystywane do pobrania danych wejściowych z pamięci globalnej i zapisania wyników z powrotem.

Planowanie warp i obsługa rozbieżności

EfTutaj znajduje się polski przekład pliku Markdown. Dla kodu, nie tłumaczono kodu, tylko komentarze.

Optymalna planowanie splotów jest kluczowe dla maksymalizacji wydajności SM. Planista splotów odpowiada za wybór splotów gotowych do wykonania i wysyłanie ich do dostępnych jednostek wykonawczych. Głównym celem planisty splotów jest utrzymanie zajętości jednostek wykonawczych przez zapewnienie, że zawsze dostępne są sploty gotowe do wykonania.

SM wykorzystuje dwupoziomowy mechanizm planowania splotów:

Planowanie splotów: Planista splotów wybiera sploty gotowe do wykonania w oparciu o strategię planowania, taką jak round-robin lub najstarszy-pierwszy. Wybrane sploty są następnie wysyłane do dostępnych jednostek wykonawczych.
Planowanie instrukcji: W obrębie każdego splotu, SM planuje instrukcje w oparciu o ich zależności i dostępność jednostek wykonawczych. SM może wydawać wiele niezależnych instrukcji z tego samego splotu w pojedynczym cyklu, aby zmaksymalizować równoległość na poziomie instrukcji.

Rysunek 7.3 ilustruje dwupoziomowy mechanizm planowania splotów.

    Pula splotów
    Splot 1 (Gotowy)
    Splot 2 (Oczekujący)
    Splot 3 (Gotowy)
    ...
    Splot N (Gotowy)
        |
        v
    Planista splotów
        |
        v
    Jednostka wysyłania
        |
        v
    Jednostki wykonawcze

Rysunek 7.3: Dwupoziomowy mechanizm planowania splotów.

Jednym z kluczowych wyzwań w planowaniu splotów jest obsługa rozbieżności gałęzi. W modelu wykonania SIMT, wszystkie wątki w splocie wykonują tę samą instrukcję w tym samym czasie. Jednak, gdy splot natrafi na instrukcję rozgałęzienia (np. instrukcja if-else), niektóre wątki mogą wykonać gałąź if, a inne gałąź else. Ta sytuacja nazywana jest rozbieżnością gałęzi.

Aby obsłużyć rozbieżność gałęzi, SM stosuje technikę nazywaną predykacją. Gdy splot napotka rozbieżną gałąź, SM wykonuje obie ścieżki gałęzi sekwencyjnie, maskując wątki, które nie biorą udziału w każdej ścieżce. Wyniki są następnie łączone za pomocą rejestrów predykatów, aby zapewnić, że każdy wątek otrzymuje prawidłowy wynik.

Przykład 7.2 przedstawia jądro CUDA z rozbieżną gałęzią.Poniżej znajduje się tłumaczenie pliku na język polski. Komentarze do kodu zostały przetłumaczone, a sam kod pozostał niezmieniony.

__global__ void divergentKernel(int *data, int *result) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    // Jeśli data[tid] jest większe od 0
    if (data[tid] > 0) {
        // Zapisz wynik jako data[tid] * 2
        result[tid] = data[tid] * 2;
    } else {
        // W przeciwnym razie zapisz wynik jako data[tid] * 3
        result[tid] = data[tid] * 3;
    }
}

Przykład 7.2: Rdzeń CUDA z rozbieżnym warunkiem.

W tym przykładzie, warunek rozgałęzienia data[tid] > 0 może spowodować, że niektóre wątki w warpie przejdą ścieżkę if, a inne ścieżkę else. SM obsługuje tę rozbieżność, wykonując obie ścieżki sekwencyjnie i maskując nieaktywne wątki w każdej ścieżce.

Rysunek 7.4 ilustruje proces predykcji dla warpów z rozbieżnymi wątkami.

    Warp (32 wątki)
    Wątek 1: data[1] = 5, result[1] = 10
    Wątek 2: data[2] = -3, result[2] = -9
    ...
    Wątek 32: data[32] = 7, result[32] = 14

    Rozbieżny warunek:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Predykcja:
    Krok 1: Wykonaj ścieżkę if z maską
        Wątek 1: result[1] = 10
        Wątek 2: (zamaskowany)
        ...
        Wątek 32: result[32] = 14

    Krok 2: Wykonaj ścieżkę else z maską
        Wątek 1: (zamaskowany)
        Wątek 2: result[2] = -9
        ...
        Wątek 32: (zamaskowany)

    Wynik końcowy:
    Wątek 1: result[1] = 10
    Wątek 2: result[2] = -9
    ...
    Wątek 32: result[32] = 14

Rysunek 7.4: Proces predykcji dla warpów z rozbieżnymi wątkami.

Używając predykcji, SM może obsługiwać rozbieżność gałęzi bez konieczności używania jawnych instrukcji rozgałęzienia lub rozbieżności przepływu sterowania. Jednak rozbieżne gałęzie mogą nadal mieć wpływ na wydajność, ponieważ SM musi wykonać obie ścieżki sekwencyjnie, zmniejszając efektywny poziom równoległości.

Rejestr i zbieracze operandów

Rejestr jest kluczowym elementem SM, zapewniając szybki dostęp do prywatnych rejestrów wątków. Każdy SM ma duży rejestr, aby wspierać wiele aktywnych wątków i umożliwić efektywne przełączanie kontekstu między warpami.Oto polski przekład tego pliku Markdown. Dla kodu, nie tłumacze kodu, tylko komentarze.

W architekturze NVIDIA Volta, każdy SM ma 64 KB rejestru, zorganizowane w 32 banki po 2 KB każdy. Rejestr jest zaprojektowany, aby zapewnić wysoką przepustowość i niską latencję dostępu w celu obsługi dużej liczby współbieżnych wątków.

Aby zminimalizować konflikty banków i poprawić wydajność, SM wykorzystuje technikę zwaną "operand collection". Operand collectors są wyspecjalizowanymi jednostkami, które gromadzą operandy z banków rejestru i dostarczają je do jednostek wykonawczych. Dzięki wykorzystaniu operand collectors, SM może zmniejszyć wpływ konfliktów banków i poprawić wykorzystanie jednostek wykonawczych.

Rysunek 7.5 przedstawia uproszczony diagram rejestru i operand collectors w SM.

    Rejestr (64 KB)
    Bank 1 (2 KB)
    Bank 2 (2 KB)
    ...
    Bank 32 (2 KB)
        |
        v
    Operand Collectors
        |
        v
    Jednostki Wykonawcze

Rysunek 7.5: Rejestr i operand collectors w SM.

Operand collectors działają poprzez gromadzenie operandów z wielu instrukcji i wielu wariantów, pozwalając SM na wydawanie instrukcji z różnych wariantów do jednostek wykonawczych w jednym cyklu. Pomaga to ukryć opóźnienie dostępu do rejestru i poprawia ogólną przepustowość SM.

Przykład 7.3 pokazuje jądro CUDA, które wykonuje iloczyn skalarny dwóch wektorów.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Każdy wątek oblicza częściową sumę iloczynu skalarnego przy użyciu przypisanego indeksu
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    // Zredukuj częściowe sumy do jednej sumy
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    // Zapisz końcową sumę
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

W tym przykładzie, każdy wątek oblicza częściową sumę iloczynu skalarnego przy użyciu przypisanego indeksu.Poniżej znajduje się tłumaczenie na język polski:

Elementy z wektorów wejściowych. Częściowe sumy są przechowywane w tablicy pamięci współdzielonej partialSum. Po obliczeniu przez wszystkie wątki ich częściowych sum, przeprowadzana jest równoległa redukcja w celu zsumowania częściowych sum i uzyskania wyniku końcowego iloczynu skalarnego.

Kolektor operandów odgrywa kluczową rolę w tym przykładzie, efektywnie gromadząc operandy do dostępów do pamięci współdzielonej i operacji arytmetycznych. Pomaga to unikać konfliktów banków i poprawia wykorzystanie jednostek wykonawczych.

Wniosek

Wieloprocesorowy strumień (streaming multiprocessor) jest podstawową jednostką obliczeniową w nowoczesnych architekturach GPU. Jego projekt koncentruje się na maksymalizacji przepustowości i ukrywaniu opóźnień pamięci poprzez kombinację drobnoziarnistego wielowątkowości, wykonywania SIMT i wydajnego zbierania operandów.

Kluczowe komponenty SM obejmują planista warp, który wybiera warpy do wykonania; stos SIMT, który obsługuje rozbieżność i zbieżność gałęzi; rejestr i kolektory operandów, które zapewniają szybki dostęp do rejestrów prywatnych wątków; oraz pamięć współdzielona i pamięć podręczna L1, które umożliwiają udostępnianie i ponowne wykorzystanie danych o małym opóźnieniu.

Ponieważ architektury GPU nadal się rozwijają, badania w obszarach takich jak obsługa rozbieżności gałęzi, planowanie warp i projekt rejestru będą kluczowe dla poprawy wydajności i efektywności przyszłych GPU. Nowe techniki, takie jak dynamiczne formowanie warp, kompaktowanie bloków wątków i pamięci podręczne ponownego wykorzystania operandów, mają potencjał, aby znacznie poprawić możliwości SM i umożliwić nowe poziomy wydajności w obciążeniach obliczeniowych równoległych.

Chapter 6 Gpu Performance Metrics and Analysis Chapter 8 Interconnect and on Chip Networks