Rozdział 3: Modele programowania równoległego w projektowaniu GPU

Jednostki przetwarzania grafiki (GPU) ewoluowały od stałofunkcyjnych akceleratorów graficznych do wysoce równoległych, programowalnych silników obliczeniowych, zdolnych do przyspieszania szerokiego spektrum aplikacji. Aby umożliwić programistom efektywne wykorzystanie masywnego paralelizmu w GPU, opracowano kilka modeli programowania równoległego i interfejsów API, takich jak NVIDIA CUDA, OpenCL i DirectCompute. Te modele programowania dostarczają abstrakcji, które pozwalają programistom wyrażać paralelizm w swoich aplikacjach, jednocześnie ukrywając niskopoziomowe szczegóły sprzętu GPU.

W tym rozdziale będziemy badać kluczowe koncepcje i zasady stojące za modelami programowania równoległego dla GPU, koncentrując się na modelu wykonania SIMT (Single Instruction, Multiple Thread), modelu programowania CUDA i interfejsach API, a także na ramach OpenCL. Omówimy również techniki mapowania algorytmów na architektury GPU w celu osiągnięcia wysokiej wydajności i efektywności.

Model wykonania SIMT (Single Instruction, Multiple Thread)

Model wykonania SIMT jest podstawowym paradygmatem używanym przez nowoczesne GPU do osiągnięcia masywnego paralelizmu. W modelu SIMT duża liczba wątków wykonuje ten sam program (nazywany jądrem) równolegle, ale każdy wątek ma własny licznik programu i może podążać za różnymi ścieżkami wykonania w zależności od jego identyfikatora wątku i danych, na których operuje.

Jądra i hierarchia wątków

Jądro GPU to funkcja, która jest wykonywana równolegle przez dużą liczbę wątków. Podczas uruchamiania jądra programista określa liczbę wątków, które mają zostać utworzone, oraz sposób, w jaki są one zorganizowane w hierarchię siatek, bloków (lub tablic współbieżnych wątków - CTA) i pojedynczych wątków.

Siatka reprezentuje całą przestrzeń problemu i składa się z jednego lub więcej bloków.
Blok jest grupą wątków, które mogą współpracować i synchronizować się ze sobą za pośrednictwem pamięci współdzielonej i barier. Wątki w obrębie bloku są wykonywane na tym samym rdzeniu GPU (nazywanym wieloprocesorowym strumieniowym).Proszę o przekład tego pliku na język polski. W przypadku kodu, nie tłumaczę kodu, tylko komentarze.

Oto plik:

Każdy wątek ma unikalne ID w obrębie swojego bloku i siatki, które można wykorzystać do obliczenia adresów pamięci i podejmowania decyzji dotyczących kontroli przepływu.

Ta hierarchiczna organizacja pozwala programistom wyrażać zarówno równoległość danych (gdzie ta sama operacja jest stosowana do wielu elementów danych), jak i równoległość zadań (gdzie różne zadania są wykonywane równolegle).

Rysunek 3.1 ilustruje hierarchię wątków w modelu wykonawczym SIMT.

            Siatka
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | Blok |
    |   |   |   |
  Wątek Wątek ...

Rysunek 3.1: Hierarchia wątków w modelu wykonawczym SIMT.

Wykonywanie SIMT

W modelu wykonawczym SIMT każdy wątek wykonuje tę samą instrukcję, ale operuje na różnych danych. Jednak, w przeciwieństwie do SIMD (Pojedyncza Instrukcja, Wiele Danych), gdzie wszystkie elementy przetwarzające wykonują się w tym samym czasie, SIMT pozwala wątkom mieć niezależne ścieżki wykonywania i dywergować na instrukcjach rozgałęzień.

Gdy warp (grupa 32 wątków w GPU NVIDIA lub 64 wątków w GPU AMD) napotyka instrukcję rozgałęzienia, sprzęt GPU ocenia warunek rozgałęzienia dla każdego wątku w warpsie. Jeśli wszystkie wątki podążają tą samą ścieżką (zbiega się), warp kontynuuje wykonywanie normalnie. Jednak, jeśli niektóre wątki podążają różnymi ścieżkami (dywergują), warp jest dzielony na dwa lub więcej podwarpów, z których każdy podąża za inną ścieżką. Sprzęt GPU szereguje wykonywanie rozbieżnych ścieżek, maskując nieaktywne wątki w każdym podwarpie. Gdy wszystkie ścieżki zostaną ukończone, podwarpy ponownie się zbiegają i kontynuują wykonywanie w tym samym czasie.

Rysunek 3.2 ilustruje wykonywanie SIMT z rozbieżnym przepływem sterowania.

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | Rozgałęzienie |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
```Poprawienie konwergencji

Rysunek 3.2: Wykonywanie SIMT z rozbieżnym przepływem sterowania.

Ten mechanizm obsługi rozbieżności pozwala SIMT na obsługę bardziej elastycznego przepływu sterowania niż SIMD, ale wiąże się to z kosztem zmniejszonej wydajności SIMD, gdy występuje rozbieżność. Programiści powinni dążyć do minimalizacji rozbieżności w obrębie warpu, aby osiągnąć optymalną wydajność.

Hierarchia pamięci

Karty GPU mają złożoną hierarchię pamięci, aby wspierać wysoką przepustowość i małe opóźnienia wymagane przez obciążenia equal. Hierarchia pamięci zwykle składa się z:

Pamięć globalna: Największa, ale najwolniejsza przestrzeń pamięci, dostępna dla wszystkich wątków w jądrze. Pamięć globalna jest zazwyczaj zaimplementowana przy użyciu pamięci o wysokiej przepustowości GDDR lub HBM.
Pamięć współdzielona: Szybka, zintegrowana pamięć na układzie, współdzielona przez wszystkie wątki w bloku. Pamięć współdzielona jest używana do komunikacji między wątkami i udostępniania danych w obrębie bloku.
Pamięć stała: Przestrzeń pamięci tylko do odczytu używana do rozsyłania danych tylko do odczytu do wszystkich wątków.
Pamięć tekstur: Przestrzeń pamięci tylko do odczytu zoptymalizowana pod kątem lokalności i dostępowana za pośrednictwem pamięci podręcznej tekstur. Pamięć tekstur jest częściej używana w obciążeniach graficznych.
Pamięć lokalna: Prywatna przestrzeń pamięci dla każdego wątku, używana do wymiany rejestrów i dużych struktur danych. Pamięć lokalna jest zwykle mapowana na pamięć globalną.

Efektywne wykorzystanie hierarchii pamięci ma kluczowe znaczenie dla osiągnięcia wysokiej wydajności na kartach GPU. Programiści powinni dążyć do maksymalizacji wykorzystania pamięci współdzielonej i minimalizacji dostępów do pamięci globalnej, aby zmniejszyć opóźnienia pamięci i wąskie gardła przepustowości.

Rysunek 3.3 ilustruje hierarchię pamięci GPU.


|            |
|   Wspólna  |
|   Pamięć   |
 ____________
      |
 ____________ 
|            |
|   Lokalna  |
|   Pamięć   |
 ____________

Rysunek 3.3: Hierarchia pamięci GPU.

Model i API programowania CUDA

CUDA (Compute Unified Device Architecture) to platforma obliczeń równoległych i model programowania opracowany przez NVIDIA do ogólnego przeznaczenia obliczeń na GPU. CUDA dostarcza zestaw rozszerzeń do standardowych języków programowania, takich jak C, C++ i Fortran, które umożliwiają programistom wyrażanie równoległości i wykorzystywanie mocy obliczeniowej kart graficznych NVIDIA.

Model programowania CUDA

Model programowania CUDA opiera się na koncepcji jąder, które są funkcjami wykonywanymi równolegle przez dużą liczbę wątków na GPU. Programista określa liczbę uruchamianych wątków i ich organizację w siatkę bloków wątków.

CUDA wprowadza kilka kluczowych abstakcji ułatwiających programowanie równoległe:

Wątek: podstawowa jednostka wykonywania w CUDA. Każdy wątek ma własny licznik programu, rejestry i lokalną pamięć.
Blok: grupa wątków, które mogą współpracować i synchronizować się ze sobą. Wątki w obrębie bloku są wykonywane na tym samym multiprocesorze strumieniowym i mogą się komunikować poprzez pamięć współdzieloną.
Siatka: zbiór bloków wątków, które wykonują to samo jądro. Siatka reprezentuje całą przestrzeń problemu i może być jednowymiarowa, dwuwymiarowa lub trójwymiarowa.

CUDA dostarcza również wbudowane zmienne (np. threadIdx, blockIdx, blockDim, gridDim), które pozwalają wątkom identyfikować siebie i obliczać adresy pamięci w oparciu o ich położenie w hierarchii wątków.

Rysunek 3.4 ilustruje model programowania CUDA.

            Siatka
       ________________
      /   /   /   /   /
     /   /   /   /   / 
    /   /   /   /   /
   /   /   /   /   /
  /__/__/__/__/__/
    |   |   |   |
    |   | Blok |
    |   |   |   |
  Wątek Wątek ...

Rysunek 3.4: Model programowania CUDA.

Hierarchia pamięci CUDA


CUDA udostępnia programiście hierarchię pamięci GPU, umożliwiając jawną kontrolę nad umieszczaniem i przemieszczaniem danych. Główne przestrzenie pamięci w CUDA to:

- Pamięć globalna: Dostępna dla wszystkich wątków w jądrze i utrzymywana pomiędzy uruchomieniami jąder. Pamięć globalna ma najwyższe opóźnienie i jest zwykle używana dla dużych struktur danych.
- Pamięć współdzielona: Szybka, wbudowana na chipie pamięć współdzielona przez wszystkie wątki w bloku. Pamięć współdzielona jest używana do komunikacji między wątkami i współdzielenia danych w obrębie bloku.
- Pamięć stała: Przestrzeń pamięci tylko do odczytu używana do rozgłaszania danych tylko do odczytu do wszystkich wątków. Pamięć stała jest buforowana i zapewnia dostęp o niskim opóźnieniu.
- Pamięć teksturowa: Przestrzeń pamięci tylko do odczytu zoptymalizowana pod kątem lokalności przestrzennej i dostępna za pośrednictwem pamięci podręcznych tekstur. Pamięć teksturowa jest częściej używana w obciążeniach graficznych.
- Pamięć lokalna: Prywatna przestrzeń pamięci dla każdego wątku, używana do przepełnienia rejestrów i dużych struktur danych. Pamięć lokalna jest zazwyczaj mapowana na pamięć globalną.

Programiści mogą alokować i przenosić dane między pamięcią hosta (CPU) i urządzenia (GPU) przy użyciu API środowiska uruchomieniowego CUDA, takich jak cudaMalloc, cudaMemcpy i cudaFree.

Rysunek 3.5 ilustruje hierarchię pamięci CUDA.

| |

Rysunek 3.5: Hierarchia pamięci CUDA.

### Synchronizacja i koordynacja w CUDA

CUDA zapewnia prymitywy synchronizacji i koordynacji, aby umożliwić współpracę i komunikację między wątkami:

- Synchronizacja barierowa: Funkcja __syncthreads()Oto polski przekład pliku w formacie Markdown:

Funkcja s() działa jako bariera, która zapewnia, że wszystkie wątki w bloku osiągnęły ten sam punkt przed przejściem dalej.
- Operacje atomowe: CUDA obsługuje operacje atomowe (np. atomicAdd, atomicExch), które pozwalają wątkom wykonywać operacje odczyt-modyfikacja-zapis na pamięci współdzielonej lub globalnej bez zakłóceń z innych wątków.
- Prymitywy poziomu warpu: CUDA zapewnia wbudowane funkcje poziomu warpu (np. __shfl, __ballot), które umożliwiają wydajną komunikację i synchronizację w obrębie warpu.

Właściwe wykorzystanie prymitywów synchronizacji i koordynacji jest kluczowe dla tworzenia poprawnych i wydajnych programów równoległych w CUDA.

Przykład 3.1 pokazuje prosty kernel CUDA, który wykonuje dodawanie wektorów.

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

int main() {
    int *a, *b, *c;
    int n = 1024;
    
    // Alokacja pamięci na hoście
    a = (int*)malloc(n * sizeof(int));
    b = (int*)malloc(n * sizeof(int));
    c = (int*)malloc(n * sizeof(int));
    
    // Inicjalizacja wektorów wejściowych
    for (int i = 0; i < n; i++) {
        a[i] = i;
        b[i] = i * 2;
    }
    
    // Alokacja pamięci na urządzeniu
    int *d_a, *d_b, *d_c;
    cudaMalloc(&d_a, n * sizeof(int));
    cudaMalloc(&d_b, n * sizeof(int));
    cudaMalloc(&d_c, n * sizeof(int));
    
    // Kopiowanie wektorów wejściowych z hosta na urządzenie
    cudaMemcpy(d_a, a, n * sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, b, n * sizeof(int), cudaMemcpyHostToDevice);
    
    // Uruchomienie kernela
    int blockSize = 256;
    int numBlocks = (n + blockSize - 1) / blockSize;
    vectorAdd<<<numBlocks,blockSize>>>(d_a, d_b, d_c, n);
    
    // Kopiowanie wektora wyniku z urządzenia na hosta
    cudaMemcpy(c, d_c, n * sizeof(int), cudaMemcpyDeviceToHost);
    
    // Zwolnienie pamięci na urządzeniu
    cudaFree(d_a);
    cudaFree(d_b); 
    cudaFree(d_c);
    
    // Zwolnienie pamięci na hoście
    free(a); 
    free(b);
    free(c);
    
    retur
```Poniżej znajduje się tłumaczenie pliku Markdown na język polski. Komentarze w kodzie zostały przetłumaczone, a sam kod pozostał niezmieniony.

n 0;
}

Ten kod CUDA uruchamia kernel vectorAdd z numBlocks blokami i blockSize wątkami na blok. Kernel wykonuje elementarne dodawanie wektorów wejściowych a i b oraz zapisuje wynik w wektorze c. Składnia <<<...>>> służy do określenia wymiarów siatki i bloków podczas uruchamiania kernela.

Strumienie i zdarzenia CUDA

Strumienie i zdarzenia CUDA zapewniają mechanizm współbieżnego wykonywania i synchronizacji kerneli oraz operacji pamięci:

Strumienie: sekwencja operacji (uruchomienia kerneli, kopiowanie pamięci) wykonywanych w kolejności. Różne strumienie mogą być wykonywane współbieżnie, co umożliwia nakładanie się obliczeń i transferów pamięci.
Zdarzenia: znaczniki, które można wstawić do strumienia, aby zapisać ukończenie określonych operacji. Zdarzenia mogą być używane do celów synchronizacji i pomiaru czasu.

Strumienie i zdarzenia umożliwiają programistom optymalizację wydajności ich aplikacji CUDA przez nakładanie obliczeń i transferów pamięci oraz wykorzystanie pełnych możliwości sprzętu GPU.

Przykład 3.2 pokazuje użycie strumieni CUDA do nakładania wykonywania kerneli i transferów pamięci.

// Utwórz dwa strumienie
cudaStream_t stream1, stream2;
cudaStreamCreate(&stream1);
cudaStreamCreate(&stream2);
 
// Asynchronicznie skopiuj dane wejściowe na urządzenie
cudaMemcpyAsync(d_a, a, size, cudaMemcpyHostToDevice, stream1);
cudaMemcpyAsync(d_b, b, size, cudaMemcpyHostToDevice, stream2);
 
// Uruchom kernele w różnych strumienach
kernelA<<<blocks, threads, 0, stream1>>>(d_a);
kernelB<<<blocks, threads, 0, stream2>>>(d_b);
 
// Asynchronicznie skopiuj wyniki z powrotem na host
cudaMemcpyAsync(a, d_a, size, cudaMemcpyDeviceToHost, stream1);
cudaMemcpyAsync(b, d_b, size, cudaMemcpyDeviceToHost, stream2);
 
// Zsynchronizuj strumienie
cudaStreamSynchronize(stream1);
cudaStreamSynchronize(stream2);

W tym przykładzie utworzono dwa strumienie CUDA. Dane wejściowe są asynchronicznie kopiowane na urządzenie przy użyciu każdego strumienia. Następnie kernele są uruchamiane w różnych strumienach, a wyniki są asynchronicznie kopiowane z powrotem na host. Na końcu strumienie są synchronizowane.Poniżej znajduje się tłumaczenie pliku Markdown na język polski. Komentarze w kodzie zostały przetłumaczone, natomiast sam kod pozostał bez zmian.

Ramowa infrastruktura OpenCL

OpenCL (Open Computing Language) to otwarty, bezpłatny standard do programowania równoległego na heterogenicznych platformach, w tym na procesorach CPU, GPU, FPGA i innych akceleratorach. OpenCL dostarcza ujednolicony model programowania oraz zestaw interfejsów API, które pozwalają deweloperom na tworzenie przenośnego i wydajnego kodu równoległego.

Model programowania OpenCL

Model programowania OpenCL jest podobny do modelu CUDA, z kilkoma kluczowymi różnicami w terminologii i abstrakcjach:

Kernel: Funkcja wykonywana równolegle przez dużą liczbę elementów roboczych (wątków) na urządzeniu OpenCL.
Element roboczy: Podstawowa jednostka wykonywania w OpenCL, analogiczna do wątku w CUDA.
Grupa robocza: Zbiór elementów roboczych, które mogą synchronizować się i udostępniać dane za pośrednictwem pamięci lokalnej. Grupy robocze są analogiczne do bloków wątków w CUDA.
NDRange: Definiuje przestrzeń indeksową i organizację elementów roboczych dla wykonania kernela. Może być jedno-, dwu- lub trójwymiarowa.

OpenCL definiuje również hierarchiczny model pamięci podobny do CUDA:

Pamięć globalna: Dostępna dla wszystkich elementów roboczych we wszystkich grupach roboczych, analogiczna do pamięci globalnej w CUDA.
Pamięć lokalna: Współdzielona przez wszystkie elementy robocze w grupie roboczej, analogiczna do pamięci współdzielonej w CUDA.
Pamięć prywatna: Prywatna dla pojedynczego elementu roboczego, analogiczna do rejestrów w CUDA.
Pamięć stała: Pamięć tylko do odczytu, dostępna dla wszystkich elementów roboczych.

Kernele OpenCL są kompilowane w czasie wykonania przez środowisko uruchomieniowe OpenCL. Program hosta może zapytać o dostępne urządzenia OpenCL, wybrać odpowiednie urządzenie, utworzyć kontekst i zbudować kernel dla tego konkretnego urządzenia. Umożliwia to aplikacjom OpenCL wysoką przenośność na różne platformy sprzętowe.

Przykład 3.3 pokazuje kernel OpenCL, który wykonuje dodawanie wektorów, podobne do przykładu CUDA z Przykładu 3.1.

// Kernel wykonujący dodawanie wektorów
__kernel void vectorAdd(__global const int *a, __global const int *b, __global int *c, int n) {
    int i = get_global_id(0);
    if (i < n)
        c[i] = a[i] + b[i];
}
```Plik Markdown:
 
```c
__kernel void vector_add(
    __global const int *a,
    __global const int *b, 
    __global int *c, 
    int n) {
    int i = get_global_id(0);
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

Tłumaczenie komentarzy:

Słowo kluczowe `__kernel` definiuje funkcję jądra OpenCL. Słowo kluczowe `__global` określa, że wskaźnik wskazuje na pamięć globalną. Funkcja `get_global_id` zwraca globalny indeks bieżącego elementu roboczego, który jest używany do obliczenia adresów pamięci dla wektorów wejściowych i wyjściowych.

### Mapowanie algorytmów na architektury GPU

Efektywne mapowanie algorytmów na architekturę GPU ma kluczowe znaczenie dla uzyskania wysokiej wydajności. Kluczowe zagadnienia obejmują:

- Ujawnianie wystarczającej równoległości: Algorytm powinien być podzielony na wiele drobnych wątków, które mogą być wykonywane współbieżnie, aby w pełni wykorzystać możliwości równoległego przetwarzania GPU.

- Minimalizacja rozbieżności gałęzi: Rozbieżny przepływ sterowania w obrębie warpu/fali może prowadzić do serializacji i zmniejszenia wydajności SIMD. Algorytmy powinny być strukturyzowane tak, aby w miarę możliwości minimalizować rozbieżność gałęzi.

- Wykorzystanie hierarchii pamięci: Dostęp do pamięci globalnej jest kosztowny. Algorytmy powinny maksymalizować wykorzystanie pamięci współdzielonej i rejestrów, aby zmniejszyć dostępy do pamięci globalnej. Dane powinny być również układane w pamięci w sposób umożliwiający scalone dostępy do pamięci.

- Bilansowanie obliczeń i dostępów do pamięci: Algorytmy powinny mieć wysoką proporcję operacji arytmetycznych do operacji pamięci, aby skutecznie ukrywać opóźnienia dostępu do pamięci i osiągać wysoką przepustowość obliczeniową.

- Minimalizacja transferów danych host-device: Przenoszenie danych między pamięcią hosta a pamięcią urządzenia jest wolne. Algorytmy powinny minimalizować takie transfery, wykonując jak najwięcej obliczeń na GPU.

Kilka wzorców projektowych równoległych algorytmów jest powszechnie używanych przy tworzeniu jąder GPU:

- Mapowanie: Każdy wątek wykonuje tę samą operację na różnych danych, umożliwiając proste przetwarzanie równoległe dużych zbiorów danych.

- Redukcja: Równoległa redukcja służy do efektywnego obliczania pojedynczej wartości (np. sumy, maksimum) z dużego zbioru danych wejściowych.
```Here is the Polish translation of the markdown file, with the code comments translated:

Wątki wykonują lokalne redukcje, które są następnie łączone w celu uzyskania ostatecznego wyniku.

- Skan: Znany również jako suma prefiksowa, skan służy do obliczania bieżącej sumy elementów w tablicy. Wydajne równoległe algorytmy skanowania są kluczowymi budulcami wielu aplikacji akcelerowanych przez GPU.

- Stencil: Każdy wątek oblicza wartość na podstawie sąsiednich elementów danych. Obliczenia stencylowe są powszechne w symulacjach naukowych i aplikacjach do obróbki obrazu.

- Zbieranie/Rozpraszanie: Wątki odczytują (zbieranie) lub zapisują (rozpraszanie) do dowolnych lokalizacji w pamięci globalnej. Staranny układ danych i wzorce dostępu są wymagane dla wydajności.

## Wniosek

Modele programowania GPU, takie jak CUDA i OpenCL, ujawniają równoległe możliwości przetwarzania nowoczesnych kart GPU programistom, umożliwiając im przyspieszenie szerokiego zakresu aplikacji. Te modele programowania dostarczają abstrakcji, które pozwalają na efektywne mapowanie drobnoziarnistych obciążeń równoległych na sprzęt GPU.

Zrozumienie modelu wykonania, hierarchii pamięci i prymitywów synchronizacji dostarczanych przez te modele programowania jest kluczowe dla tworzenia wydajnego kodu GPU. Programiści muszą starannie rozważać takie czynniki, jak organizacja wątków, rozbieżność gałęzi, wzorce dostępu do pamięci i projektowanie algorytmów, aby w pełni wykorzystać moc obliczeniową kart GPU.

Ponieważ architektury GPU nadal ewoluują, modele programowania i narzędzia również muszą się rozwijać, aby umożliwić programistom efektywne wykorzystanie nowych funkcji i możliwości sprzętu. Trwające badania w obszarach, takich jak projektowanie języków programowania, optymalizacja kompilatora i autotuning, będą kluczowe dla poprawy wydajności projektowania i przenośności w erze obliczeniowej heterogenicznej.

Chapter 2 Gpu Rogramming Models Chapter 4 Gpu Memory System Design