Rozdział 2: Modele Programowania GPU

Jednostki przetwarzania graficznego (GPU) ewoluowały od stałofunkcyjnych akceleratorów graficznych do wysoce równoległych, programowalnych silników obliczeniowych, zdolnych do przyspieszania szerokiego zakresu aplikacji. Aby umożliwić programistom efektywne wykorzystanie masywnej równoległości w GPU, opracowano kilka modeli programowania równoległego i interfejsów API, takich jak NVIDIA CUDA, OpenCL i DirectCompute. Te modele programowania zapewniają abstrakcje, które pozwalają programistom wyrażać równoległość w swoich aplikacjach, jednocześnie ukrywając niskie szczegóły sprzętowe GPU.

W tym rozdziale zbadamy kluczowe koncepcje i zasady stojące za modelami programowania równoległego dla GPU, koncentrując się na modelu wykonawczym, architekturach zestawu instrukcji (ISA) GPU, architekturze ISA NVIDIA GPU oraz architekturze ISA AMD Graphics Core Next (GCN). Podamy także przykłady ilustrujące, jak te koncepcje są stosowane w praktyce.

Model Wykonawczy

Model wykonawczy nowoczesnych modeli programowania GPU opiera się na pojęciu jąder, które są funkcjami wykonywanymi równolegle przez dużą liczbę wątków na GPU. Przy uruchamianiu jądra, programista określa liczbę wątków do utworzenia i w jaki sposób są one zorganizowane w hierarchię siatek, bloków (lub tablic współpracujących wątków - CTAs) oraz 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 ramach bloku są wykonywane na tym samym rdzeniu GPU (nazywanym multiprocesorowym strumieniem lub jednostką obliczeniową).
Każdy wątek ma unikatowy identyfikator w obrębie swojego bloku i siatki, który może być używany do obliczania adresów pamięci i podejmowania decyzji dotyczących przepływu sterowania.

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).

RysunekOto tłumaczenie na język polski:

e 2.1 ilustruje hierarchię wątków w modelu wykonywania na GPU.

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

Rysunek 2.1: Hierarchia wątków w modelu wykonywania na GPU.

Wykonywanie SIMT

Modele programowania GPU, takie jak CUDA i OpenCL, stosują model wykonywania Single-Instruction, Multiple-Thread (SIMT). W modelu SIMT, wątki są wykonywane w grupach nazywanych warpsami (terminologia NVIDIA) lub wavefrontami (terminologia AMD). Wszystkie wątki w obrębie warpsa wykonują tę samą instrukcję w tym samym czasie, ale każdy wątek działa na różnych danych.

Jednak, w odróżnieniu od tradycyjnego modelu Single-Instruction, Multiple-Data (SIMD), gdzie wszystkie elementy przetwarzające wykonują się w synchronizacji, SIMT pozwala wątkom mieć niezależne ścieżki wykonywania i rozbiegać się na instrukcjach warunkowych. Gdy warp napotyka instrukcję warunkową, sprzęt GPU ocenia warunek rozgałęzienia dla każdego wątku w warpie. Jeśli wszystkie wątki podążają tą samą ścieżką (zbieżne), warp kontynuuje wykonywanie normalnie. Jeśli niektóre wątki podążają różnymi ścieżkami (rozbieżne), warp jest dzielony na dwa lub więcej podwarpsó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 synchronizacji.

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

         Warp
   ________________
  /   /   /   /   /
 /   /   /   /   / 
/   /   /   /   /
   |   |   |
   | Rozgałęzienie |
   |   |   |
  / \ / \ / \
 /   X   \   \
/   / \   \   \
   /   \   \
  /     \   \
 /       \   \
/         \   \
           \
            \
             \
   Ponowne zbieżenie

Rysunek 2.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.Proszę o tłumaczenie tego pliku Markdown na język polski. Dla kodu, nie tłumacz kodu, tylko tłumacz komentarze.

Hierarchia pamięci

Karty graficzne mają złożoną hierarchię pamięci, aby wspierać wysoką przepustowość i niskie opóźnienia wymagane przez równoległe obciążenia robocze. 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 zwykle zaimplementowana przy użyciu pamięci GDDR lub HBM o dużej przepustowości.
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 udostępniania 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ęć tekstur: Przestrzeń pamięci tylko do odczytu zoptymalizowana pod kątem lokalności przestrzennej i dostępna 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 przepełnienia rejestrów i dużych struktur danych. Pamięć lokalna jest zwykle mapowana na pamięć globalną.

Skuteczne wykorzystanie hierarchii pamięci jest kluczowe dla osiągnięcia wysokiej wydajności na kartach graficznych. 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 2.3 ilustruje hierarchię pamięci GPU.

      ____________
     |            |
     |   Pamięć   |
     |   Globalna  |
      ____________
           |
      ____________
     |            |
     |  Pamięć    |
     |   Stała    |
      ____________
           |
      ____________
     |            |
     |  Pamięć    |
     |   Tekstur   |
      ____________
           |
           |
      ____________
     |            |
     |  Pamięć    |
     |  Współdzielona|
      ____________
           |
      ____________ 
     |            |
     |  Pamięć    |
     |   Lokalna   |
      ____________

Rys.Oto tłumaczenie pliku Markdown na język polski. Komentarze w kodzie zostały przetłumaczone, a sam kod pozostał niezmieniony.

Architektury zestawu instrukcji GPU

Architektury zestawu instrukcji GPU (ISA) definiują niskopoziomowy interfejs między oprogramowaniem a sprzętem. Określają one instrukcje, rejestry i sposoby adresowania pamięci obsługiwane przez GPU. Zrozumienie architektur ISA GPU jest kluczowe dla opracowywania wydajnego kodu GPU i optymalizacji wydajności.

W tej sekcji zbadamy architektury ISA dwóch głównych dostawców GPU: NVIDII i AMD. Skupimy się na architekturach Parallel Thread Execution (PTX) i SASS NVIDII oraz architekturze Graphics Core Next (GCN) AMD.

Architektury ISA GPU NVIDII

Karty graficzne NVIDII obsługują dwa poziomy architektur ISA: PTX (Parallel Thread Execution) i SASS (Streaming ASSembler). PTX to wirtualna architektura ISA, która zapewnia stabilny cel dla kompilatorów CUDA, podczas gdy SASS jest natywną architekturą ISA kart graficznych NVIDII.

PTX (Parallel Thread Execution)

PTX to niskopoziomowa, wirtualna architektura ISA zaprojektowana dla kart graficznych NVIDII. Jest ona podobna do LLVM IR lub bajtkodu Java, ponieważ zapewnia stabilny, niezależny od architektury cel dla kompilatorów. Programy CUDA są zwykle kompilowane do kodu PTX, który jest następnie tłumaczony na natywne instrukcje SASS przez sterownik karty graficznej NVIDII.

PTX obsługuje szeroką gamę instrukcji arytmetycznych, pamięciowych i sterowania przepływem. Posiada nieograniczoną liczbę wirtualnych rejestrów i obsługuje predykację, co umożliwia wydajną implementację sterowania przepływem. PTX zapewnia również specjalne instrukcje do synchronizacji wątków, operacji atomowych i pobierania próbek tekstur.

Oto przykład kodu PTX dla prostego jądra dodawania wektorów:

.version 7.0
.target sm_70
.address_size 64

.visible .entry vecAdd(
    .param .u64 vecAdd_param_0,
    .param .u64 vecAdd_param_1,
    .param .u64 vecAdd_param_2,
    .param .u32 vecAdd_param_3
)
{
    .reg .b32 %r<4>;
    .reg .b64 %rd<6>;

    // Wczytaj parametry z pamięci
    ld.param.u64 %rd1, [vecAdd_param_0];
    ld.param.u64 %rd2, [vecAdd_param_1];
    ld.param.u64 %rd3, [vecAdd_param_2];
    ld.param.u32 %r1, [vecAdd_param_3];
    cvta.to.global.u64 %rd4, %rd1;
    cvta
```Oto tłumaczenie pliku Markdown na język polski. Komentarze w kodzie zostały przetłumaczone, natomiast sam kod nie został przetłumaczony.

// Przekonwertuj do globalnego adresu 64-bitowego .to.global.u64 %rd5, %rd2; // Pobierz ID wątku globalnego mov.u32 %r2, %tid.x; // Pomnóż ID wątku przez rozmiar jednego elementu (4 bajty) mul.wide.u32 %rd6, %r2, 4; // Dodaj przesunięcie do wskaźnika na pierwszy element wejściowego wektora add.s64 %rd7, %rd4, %rd6; // Dodaj przesunięcie do wskaźnika na pierwszy element wyjściowego wektora add.s64 %rd8, %rd5, %rd6;

// Załaduj element wejściowy z pamięci globalnej ld.global.f32 %f1, [%rd7]; // Załaduj element wejściowy z pamięci globalnej ld.global.f32 %f2, [%rd8]; // Dodaj elementy wejściowe add.f32 %f3, %f1, %f2;

// Przekonwertuj do globalnego adresu 64-bitowego cvta.to.global.u64 %rd9, %rd3; // Oblicz adres elementu wyjściowego add.s64 %rd10, %rd9, %rd6; // Zapisz wynik do pamięci globalnej st.global.f32 [%rd10], %f3;

// Zwróć z funkcji ret; }


To funkcja kernela `vecAdd`, która przyjmuje cztery parametry: wskaźniki na wektory wejściowe i wyjściowy oraz rozmiar wektorów. Kernel oblicza globalny identyfikator wątku, ładuje odpowiednie elementy z wektorów wejściowych, wykonuje dodawanie i zapisuje wynik do wektora wyjściowego.

#### SASS (Streaming ASSembler)

SASS to natywny ISA (Instruction Set Architecture) kart graficznych NVIDIA. Jest to niskopoziomowy, specyficzny dla sprzętu ISA, który bezpośrednio mapuje się na sprzęt GPU. Instrukcje SASS są generowane przez sterownik GPU NVIDIA z kodu PTX i zazwyczaj nie są widoczne dla programistów.

Instrukcje SASS są kodowane w kompaktowym formacie, aby zmniejszyć zapotrzebowanie na przepustowość pamięci i pojemność pamięci podręcznej instrukcji. Obsługują one szeroki zakres typów operandów, w tym rejestry, wartości bezpośrednie i różne tryby adresowania dla dostępu do pamięci.

Oto przykład kodu SASS dla jądra dodawania wektorów:

```sass
code_version_number 90
                     @P0 LDG.E R2, [R8];
                     @P0 LDG.E R0, [R10];
                     @P0 FADD R0, R2, R0;
                     @P0 STG.E [R12], R0;
                         EXIT;

Ten kod SASS odpowiada wcześniejszemu kodowi PTX. Ładuje on elementy wektorów wejściowych z pamięci globalnej (LDG.E), wykonuje dodawanie (FADD), zapisuje wynik z powrotem do pamięci globalnej (STG.E) i kończy działanie jądra.

Architektura i ISA AMD Graphics Core Next

Karty graficzne AMD wykorzystują architekturę i ISA Graphics Core Next (GCN). GCN jest opartym na RISC ISA, które obsługuje zarówno grafiki, jak i obliczenia. Zaprojektowano go z myślą o wysokiej wydajności, skalowalności i efektywności energetycznej.

GCN wprowadza kilka kluczowych funkcji, takich jak:

Architektura RISC ukierunkowana na obliczenia
Wielokrotne jednostki wykonawcze (wavefronts) w celu zwiększenia wydajności
Efektywne wykorzystanie pamięci i zasobów GPU
Obsługa wielu zestawów instrukcji, w tym 32-bitowych i 64-bitowych
Zoptymalizowana pod kątem wysokiej wydajności i energooszczędnościHere is the Polish translation of the provided markdown file, with the code comments translated:

alar ALU do wydajnego wykonywania operacji skalarnych i sterowania przepływem.

Wektor ALU do równoległego wykonywania operacji równoległych danych.
Wysoko wydajny system pamięci z obsługą operacji atomowych i dostępu o niskim opóźnieniu do pamięci współdzielonej.
Elastyczny tryb adresowania dla operacji pamięci, obsługujący adresowanie base+offset oraz skalar+wektor.

Oto przykład kodu ISA GCN dla jądra dodawania wektorów:

.text
.globl vecAdd
.p2align 2
 
.type vecAdd,@function
vecAdd:
    .set DPTR, 0
 
    # Załaduj argumenty jądra z pamięci
    s_load_dwordx4 s[0:3], s[4:5], 0x0
    s_load_dword s4, s[4:5], 0x10
    s_waitcnt lgkmcnt(0)
 
    # Przygotuj adresy pamięci
    v_lshlrev_b32 v0, 2, v0
    v_add_u32 v1, vcc, s1, v0
    v_mov_b32 v3, s3
    v_addc_u32 v2, vcc, s2, v3, vcc
    flat_load_dword v1, v[1:2]
 
    v_add_u32 v3, vcc, s0, v0
    v_mov_b32 v5, s3
    v_addc_u32 v4, vcc, s2, v5, vcc
    flat_load_dword v0, v[3:4]
 
    # Wykonaj dodawanie wektorów
    v_add_f32 v0, v0, v1
    flat_store_dword v[3:4], v0
    s_endpgm

W tym kodzie GCN ładowane są elementy wejściowych wektorów przy użyciu flat_load_dword, wykonywane jest dodawanie za pomocą v_add_f32 i zapisywany jest wynik do pamięci przy użyciu flat_store_dword. Polecenia s_load_dwordx4 i s_load_dword służą do załadowania argumentów jądra z pamięci.

Mapowanie algorytmów na architektury GPU

Efektywne mapowanie algorytmów na architekturę GPU jest kluczowe dla osiągnięcia wysokiej wydajności. Kluczowe kwestie to:

Ujawnianie wystarczającej ilości równoległości

Algorytm powinien być rozbity na wiele drobnych wątków, które mogą być wykonywane równolegle, aby w pełni wykorzystać równoległe możliwości przetwarzania GPU. Często oznacza to identyfikację części algorytmu, które mogą być wykonywane niezależnie na różnych elementach danych.

Minimalizacja rozbieżności gałęzi

Rozbieżny przepływ sterowania w obrębie warpu/wavefrontu może prowadzić do serializacji i zmniejszonej wydajności SIMD. Algorytmy powinny być strukturyzowane w celu zminimalizowania rozbieżności gałęzi, w miarę możliwości. Można to osiągnąć poprzez ograniczenie stosowania zależnego od danych sterowania.Oto tłumaczenie pliku na język polski. W przypadku kodu, nie tłumaczono kodu, tylko komentarze.

Wykorzystywanie 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ż ułożone w pamięci w sposób umożliwiający scalone dostępy do pamięci, gdzie wątki w warpie uzyskują dostęp do sąsiadujących lokalizacji w pamięci. Efektywne wykorzystanie hierarchii pamięci może znacząco zmniejszyć opóźnienia w dostępie do pamięci i problemy z przepustowością.

Równoważenie obliczeń i dostępów do pamięci

Algorytmy powinny mieć wysoki stosunek operacji arytmetycznych do operacji pamięci, aby skutecznie ukrywać opóźnienia w dostępie do pamięci i osiągać wysoką przepustowość obliczeniową. Można to osiągnąć przez maksymalizację ponownego wykorzystania danych, wstępne pobieranie danych i nakładanie obliczeń na dostępy do pamięci.

Minimalizowanie transferów danych między hostem a urządzeniem

Transferowanie danych między pamięcią hosta (CPU) a pamięcią urządzenia (GPU) jest wolne. Algorytmy powinny minimalizować takie transfery, wykonując jak najwięcej obliczeń na GPU. Dane powinny być przenoszone na GPU w dużych partiach i przechowywane na urządzeniu tak długo, jak to konieczne, aby amortyzować koszt transferu.

Przy opracowywaniu kerneli GPU często stosuje się kilka wzorców projektowych algorytmów równoległych:

Mapowanie: Każdy wątek wykonuje tę samą operację na innym elemencie danych, umożliwiając proste równoległe przetwarzanie dużych zbiorów danych.
Redukcja: Równoległa redukcja służy do wydajnego obliczania pojedynczej wartości (np. sumy, maksimum) z dużego zbioru danych wejściowych. Wątki wykonują lokalne redukcje, które są następnie łączone, aby uzyskać wynik końcowy.
Skan: Znany również jako suma prefiksowa, skan służy do obliczenia bieżącej sumy elementów w tablicy. Wydajne równoległe algorytmy skanowania są kluczowymi blokami konstrukcyjnymi dla wielu aplikacji GPU-akcelerowanych.
Szablony: Każdy wątek oblicza wartość na podstawie sąsiednich elementów danych. Obliczenia szablonowe są powszechne w symulacjach naukowych i przetwarzaniu obrazów.Tutaj jest polski przekład pliku Markdown. Dla kodu nie tłumaczono kodu, tylko komentarze.
Zbieranie/Rozpraszanie: Wątki czytają z (zbieranie) lub zapisują do (rozpraszanie) dowolnych lokalizacji w pamięci globalnej. Uważne rozmieszczenie danych i wzorce dostępu są wymagane dla wydajności.

Rysunek 3.20 ilustruje przykład wzorca mapowania, gdzie każdy wątek stosuje funkcję (np. pierwiastek kwadratowy) do innego elementu tablicy wejściowej.

Tablica wejściowa:  
                |  |   |   |   |   |   |   |
                v  v   v   v   v   v   v   v
               ______________________________
Wątki:        |    |    |    |    |    |    |    |
              |____|____|____|____|____|____|____|
                 |    |    |    |    |    |    |
                 v    v    v    v    v    v    v
Tablica wyjściowa:

Rysunek 3.20: Przykład wzorca mapowania w programowaniu GPU.

Wniosek

Modele programowania GPU, takie jak CUDA i OpenCL, odsłaniają możliwości przetwarzania równoległego 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 pisania wysokowydajnego kodu GPU. Programiści muszą starannie rozważyć czynniki takie 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 muszą również postępować naprzód, aby umożliwić programistom efektywne wykorzystanie nowych funkcji i możliwości sprzętu. Bieżące badania w obszarach takich jak projektowanie języków programowania, optymalizacja kompilatora i automatyczne dostrajanie będą kluczowe dla poprawy wydajności i produktywności programistów w erze obliczeń heterogenicznych.

Chapter 1 Introduction to Gpu Chip Design Chapter 3 Parallel Programming Models