CUDA - 10 tutoriali dla początkujących

Podstawy programowania równoległego z użyciem CUDA

View the Project on GitHub Kowalikov/CUDA_blogs

II. Sprawdzanie środowiska CUDA

Autor: Natan Jarzyński, Marek Kowalik

Słowa kluczowe: GPU programming, programming, coding, CUDA toolkit, CUDA, C++, GPU, programowanie równoległe, instalacja CUDA, pierwszy program CUDA, alokacja pamięci GPU, kernel CUDA, nvcc, nvidia-smi

CUDA w C++ - pierwsze testy

W drugiej odsłonie tej serii sprawdzimy, jak działa GPU z CUDA w C++.

🧰 Kompilowanie skryptów CUDA C++ '.cu' z 'nvcc'

📌 Przypominajka, jak skompilować i uruchomić program CUDA na Linuxie (na Windowsie będzie podobnie, ale z rozszerzeniem .exe):

nvcc kernel.cu -o kernel && chmod u+x ./kernel && ./kernel

🧰 Test 1: Sprawdzamy, czy CUDA widzi GPU

Bez GPU ani rusz. Najpierw sprawdzimy, czy CUDA jest w stanie wykryć naszą kartę graficzną NVIDIA. W tym celu wywołamy funkcję cudaGetDeviceCount, która zwraca liczbę dostępnych urządzeń CUDA. Jeśli wynik jest większy niż 0, oznacza to, że GPU jest dostępne i gotowe do pracy z CUDA. Alternatywnie, środowisko CUDA mogło zostać niepoprawnie skonfigurowane, i sama funkcja zwróci błąd - ten przypadek też obsługujemy.

#include <iostream>
#include <cuda_runtime.h>

    
int main() {
    int count = 0;
    if (cudaGetDeviceCount(&count) == cudaSuccess && count > 0)
        printf("CUDA device available\n");
    else
        printf("No CUDA device!\n");

    return 0;
}

🧰 Test 2: Sprawdzamy, jakie obliczenia możemy puścić na GPU

Jeżeli widzimy GPU, to sprawdźmy, jakie daje nam ona możliwości. Do zaznajomienia się ze specyfikacją GPU użyjemy funkcji cudaDeviceGetAttribute, która pozwala nam odczytać różne parametry dotyczące możliwości naszego GPU. Zwraca ona wartość artybutu, o który pytamy w drugim argumencie funkcji, pod zmienną, której adres podajemy w pierwszym argumencie. Ostatni argument to numer urządzenia (zazwyczaj 0, jeśli mamy tylko jedno GPU). Teraz, my będziemy pytać o:

Trochę tego dużo, ale efektywnie pozwoli nam to zrozumieć, jak duże obliczenia możemy uruchomić na naszym GPU i jak je najlepiej zorganizować. Rozpakujmy ten żargon.

Hierarchia wykonania CUDA — wątki, bloki, siatki i multiprocesory w pigułce

1) Wątki: podstawowy element pracy równoległej W CUDA wątek jest najmniejszą jednostką wykonawczą. Każdy wątek uruchamia tę samą funkcję jądra (kernel), ale zazwyczaj przetwarza inny element danych (np. jeden piksel, jeden indeks tablicy, jedną cząsteczkę). Wątki mają własne rejestry i prywatną pamięć lokalną, a także dostęp do pamięci współdzielonej w obrębie bloku oraz pamięci globalnej urządzenia. Na GPU NVIDII wątki są wykonywane w grupach po 32, zwanych warpami. Warp pracuje synchronicznie na jednym Streaming Multiprocessorze (SM). Zachowanie spójnego przepływu sterowania i koaleskowalnych dostępów do pamięci w obrębie warpa jest kluczowe dla wysokiej wydajności.

2) Bloki wątków: współpraca, pamięć współdzielona i synchronizacja Blok wątków to grupa wątków, które mogą ze sobą współpracować. Wątki w tym samym bloku mogą synchronizować się za pomocą __syncthreads() oraz komunikować poprzez szybką, niskolatencyjną pamięć współdzieloną. Dzięki temu bloki świetnie nadają się do takich wzorców obliczeń jak tiling, redukcje, skany czy operacje stencylowe. Bloki mogą być jednowymiarowe, dwuwymiarowe lub trójwymiarowe (w sensie zbierają w sobie wątki, które można indeksować po wymiarach X, Y, Z) — zależnie od sposobu odwzorowania danych (np. obrazów). Są też niezależne od siebie: sprzęt może wykonywać je w dowolnej kolejności i na różnych SM-ach. Z tego powodu kernel nie powinien zakładać istnienia synchronizacji pomiędzy blokami (poza oddzielnymi wywołaniami kernela lub użyciem wyspecjalizowanych API).

3) Siatki (grids): uruchamianie wielu bloków, aby pokryć cały problem Siatka (grid) to zbiór wszystkich bloków uruchomionych w pojedynczym wywołaniu kernela. Podobnie jak bloki, siatki mogą mieć 1–3 wymiary i są tak definiowane, aby pokryć cały przetwarzany zakres danych. Siatka określa całkowity poziom równoległości, jaki będzie dostępny dla GPU. Wątek oblicza swoje globalne indeksy za pomocą wbudowanych zmiennych (blockIdx, blockDim, threadIdx, gridDim). Dzięki temu, że bloki są niezależne, ten sam kernel może skalować się na różne rozmiary danych i różne modele GPU — runtime po prostu rozdziela bloki na dostępne SM-y aż do ukończenia pracy.

4) Multiprocesory strumieniowe (SM): miejsce, gdzie dzieje się wykonanie Streaming Multiprocessor (SM) to jednostka sprzętowa, która faktycznie wykonuje warpy. Każdy SM posiada rdzenie obliczeniowe, pliki rejestrów, pamięć współdzieloną/pamięć L1 oraz harmonogramy warpów. Po uruchomieniu kernela GPU mapuje bloki na dostępne SM-y; pojedynczy SM może utrzymywać wiele bloków jednocześnie, jeżeli pozwalają na to zasoby (rejestry na wątek, pamięć współdzielona na blok, maksymalna liczba aktywnych wątków/warpów/bloków). Occupancy — czyli ile warpów jest aktywnych w stosunku do maksymalnych możliwości SM — wpływa na zdolność ukrywania opóźnień pamięci. Jednak maksymalna occupancy nie zawsze oznacza maksymalną wydajność: kluczowe jest znalezienie balansu między użyciem zasobów a liczbą aktywnych warpów.

5) Relacje między nimi: planowanie, wydajność i decyzje projektowe Podsumowując: host uruchamia siatkę złożoną z bloków; GPU przypisuje te bloki do SM-ów; SM wykonuje je jako warpy zbudowane z wątków. Wynikają z tego ważne implikacje dla wydajności. Bloki powinny mieć rozmiary będące wielokrotnością 32, aby unikać częściowych warpów. Pamięci współdzielonej należy używać do ponownego wykorzystania danych i optymalizacji lokalności — z uwzględnieniem tego, że jej użycie ogranicza liczbę bloków rezydujących na SM. Dostępy do pamięci globalnej powinny być koaleskowalne w obrębie warpa. I wreszcie, logika kernela powinna opierać się na niezależności bloków, co umożliwia schedulerowi elastyczny podział pracy i naturalne skalowanie między różnymi GPU.

Puśćmy teraz kod, który odczyta te parametry i wyświetli je na ekranie:

#include <iostream>
#include <cuda_runtime.h>


void deviceSpecification(
    int& maxThreadsPerBlock,
    int& maxBlocksPerGrid,
    int& maxThreadsPerSM,
    int& multiProcessorCount
) {

    // Maximum number of threads per block
    cudaDeviceGetAttribute(&maxThreadsPerBlock, cudaDevAttrMaxThreadsPerBlock, 0);

    // Maximum number of blocks in 1D (X axis)
    cudaDeviceGetAttribute(&maxBlocksPerGrid, cudaDevAttrMaxGridDimX, 0);

    // Maximum number of threads per multiprocessor
    cudaDeviceGetAttribute(&maxThreadsPerSM, cudaDevAttrMaxThreadsPerMultiProcessor, 0);

    // Number of multiprocessors
    cudaDeviceGetAttribute(&multiProcessorCount, cudaDevAttrMultiProcessorCount, 0);

    std::cout << "*** Your GPU specifications ***\n";
    std::cout << "Max threads per block: " << maxThreadsPerBlock << "\n";
    std::cout << "Max blocks per grid (X): " << maxBlocksPerGrid << "\n";
    std::cout << "Max threads per multiprocessor: " << maxThreadsPerSM << "\n";
    std::cout << "Number of multiprocessors: " << multiProcessorCount << "\n";
}

int main() {

    int maxThreads = -1, maxBlocks = -1, maxThreadsPerSM = -1, multiProcessorCount = -1;

    // Show GPU specification
    deviceSpecification(maxThreads, maxBlocks, maxThreadsPerSM, multiProcessorCount);
    
    return 0;
}

Parę komentarzy co do tego kodu:

W moim przypadku, GPU pokazał następujące parametry:

*** Your GPU specifications ***
Max threads per block: 1024
Max blocks per grid (X): 2147483647
Max threads per multiprocessor: 1024
Number of multiprocessors: 14

To oznacza, że mogę uruchomić do 1024 wątków w jednym bloku, mogę mieć ogromną liczbę bloków w siatce (2^31-1), a na jednym multiprocesorze mogę mieć do 1024 wątków. Mój GPU ma 14 multiprocesorów, więc teoretycznie może obsłużyć do 14 * 1024 = 14336 wątków jednocześnie, ale to zależy od wielu czynników, takich jak zasoby pamięci i harmonogramowanie. Te informacje są kluczowe, gdy będziemy projektować nasze obliczenia na GPU, aby upewnić się, że wykorzystujemy jego możliwości w sposób efektywny.

Po więcej informacji o hierarchii wykonania CUDA i jej implikacjach dla projektowania kerneli, polecam zajrzeć do dokumentacji NVIDII oraz przewodnika CUDA C++ Programming Guide, gdzie jest to omówione bardzo szczegółowo.

🔍 Podsumowanie

W tym wpisie sprawdziliśmy, czy nasze środowisko CUDA jest poprawnie skonfigurowane i czy możemy wykryć nasz GPU. Następnie odczytaliśmy specyfikację naszego GPU, co pozwoli nam lepiej zrozumieć, jakie obliczenia możemy na nim uruchomić i jak je zoptymalizować. W kolejnych częściach tej serii będziemy już pisać konkretne programy CUDA, które będą wykorzystywać te możliwości do wykonywania różnych zadań obliczeniowych.

Pytania kontrolne

  1. Jakie funkcje CUDA używamy do sprawdzenia dostępności GPU?
  2. Funkcja cudaDeviceGetAttribute pozwala nam odczytać różne parametry dotyczące możliwości naszego GPU.
    (a) Jakie argumenty przyjmuje?
    (b) Jakie atrybuty znasz, które możesz odczytać?

Ćwiczenia:

  1. Sprawdź specyfikację swojego GPU i zapisz kilka kluczowych atrybutów. Puść kod z tego wpisu i porównaj wyniki z dokumentacją swojego GPU, aby lepiej zrozumieć jego możliwości.
  2. Policz ile wątków możesz uruchomić jednocześnie na swoim GPU, biorąc pod uwagę liczbę multiprocesorów i maksymalną liczbę wątków na multiprocesor. Jak to się ma do maksymalnej liczby wątków na blok?
  3. Zmodyfikuj kod, aby odczytać i wyświetlić dwa inne atrybuty GPU, które wydają Ci się interesujące, takie jak ilość ilość pamięci współdzielonej na blok, czy rozmiar warpa.

Odpowiedzi do pytań kontrolnych:

  1. Do sprawdzenia dostępności GPU używamy funkcji cudaGetDeviceCount, która odkłada liczbę dostępnych urządzeń CUDA pod przekazany wskaźnik typu int. Jeśli operacja się powiodła, funkcja zwraca cudaSuccess, i możemy odczytać ilość dostępnych GPU.
  2. (a) Funkcja cudaDeviceGetAttribute, za dokumentacją przyjmuje:
    • wskaźnik do zmiennej typu int, gdzie odłożymy wynik
    • atrybut który chcemy odczytać (typ cudaDeviceAttr)
    • numer urządzenia (zazwyczaj 0, jeśli mamy tylko jedno GPU)

    (b) Przykładowe atrybuty, które możemy odczytać to:

    • cudaDevAttrMaxThreadsPerBlock - maksymalna liczba wątków, które mogą być uruchomione w jednym bloku.
    • cudaDevAttrMaxGridDimX - maksymalna liczba bloków, które mogą być uruchomione w jednym wymiarze siatki (grid).
    • cudaDevAttrMaxThreadsPerMultiProcessor - maksymalna liczba wątków, które mogą być uruchomione na jednym multiprocesorze.
    • cudaDevAttrMultiProcessorCount - liczba multiprocesorów w GPU.
    • cudaDevAttrMaxSharedMemoryPerBlock - maksymalna ilość pamięci współdzielonej dostępna na blok.
    • cudaDevAttrWarpSize - liczba wątków w warpie (zazwyczaj 32).
    • cudaDevAttrMaxRegistersPerBlock - maksymalna liczba rejestrów dostępna na blok.

Strona główna

Poprzedni wpis: I. Ustawienie środowiska CUDA na Windowsie i Linuxie

Następny wpis: III. Pisanie funkcji na GPU