Przejdź do treści

CUDA

logo CUDA

Wprowadzenie

CUDA1 to rozwijany przez firmę NVIDIA model programowania dla platform akcelerowanych z układami GPU NVIDIA, przewidziany wyłącznie dla kart tego producenta. Standardowo umożliwia zlecanie obliczeń i komunikację pomiędzy procesorem a kartami GPU, oraz bezpośredni dostęp do pamięci pomiędzy wieloma kartami w obrębie jednego komputera. Posiada dodatkowe rozwiązania wspierające efektywną pracę z wieloma GPU, również w architekturze rozproszonej.

Dostępność

Model CUDA pierwotnie powstał dla języków C / C++, obecnie wspiera również język Fortran. Realizowany poprzez rozszerzenie składni bazowego języka oraz bibliotekę obsługującą funkcjonalności związane z GPU. Wymaga specjalnego kompilatora rozpoznającego składnię CUDA (nvcc, nvfortran). Aby korzystać z CUDA wymagane jest zainstalowanie podstawowego pakietu CUDA Toolkit lub rozszerzonego NVIDIA HPC SDK (więcej w sekcji informacje o standardzie️).

Dostępny także dla języka Python poprzez oficjalną bibliotekę pośredniczącą CUDA Python (rozwijaną dopiero od 2021 roku). W tym przypadku zwykły kod oraz wywołania API CUDA wyrażają się w języku Python, natomiast kod akcelerowany (kernele) przekazuje się jako literał zawierający fragmenty kodu w formie tekstowej, w składni CUDA C/C++.

Dodatkowe rozwiązania dla języka Python
  • PyCUDA, biblioteka podobna w założeniach do CUDA Python,
  • Numba, specjalny kompilator pozwalający wyrażać kernele w języku Python,
  • CuPy, rozbudowana biblioteka posiadająca możliwość tworzenia własnych kerneli.

Oprócz powyższych, są także dostępne nieoficjalne biblioteki pośredniczące dla innych języków.

Inne technologie

CUDA może być używana w połączeniu z innymi technologiami odpowiedzialnymi za zrównoleglanie obliczeń. Dotyczy to zarówno rozpraszania obliczeń między wiele komputerów (np. MPI) jak i zrównoleglania obliczeń wykonywanych na CPU (np. OpenMP). Dodatkowo może być łączona z innymi technikami umożliwiającymi wykorzystanie kart GPU NVIDIA, takimi jak OpenMP offloading czy OpenACC. Warto wspomnieć, że współcześnie większość głównych implementacji standardu MPI potrafi współpracować z modelem CUDA (tzw. CUDA-aware MPI), poprzez możliwość przesyłania danych bezpośrednio z pamięci GPU.

NVIDIA udostępnia również własną bibliotekę NCCL (NVIDIA Collective Communication Library), która dostarcza zbliżone do MPI (przesyłanie wiadomości, operacje kolektywne) mechanizmy wymiany danych pomiędzy wieloma GPU w modelu rozproszonym. NCCL może zostać wykorzystane zarówno do komunikacji między GPU w ramach jednego komputera, jak i na platformach w których akceleratory GPU umieszczone są w wielu węzłach obliczeniowych.

Szczegóły

Podstawową cechą programów CUDA jest możliwość podziału kodu źródłowego na część zwykłą tzw. kod hosta (wykonywaną przez CPU) oraz część akcelerowaną (wykonywaną na GPU). Rozróżnienie odbywa się tylko na poziomie specyfikatorów funkcji. Zarówno kod zwykły jak i akcelerowany implementuje się w tym samym języku (C11/C++17/Fortan 2003), z drobnymi modyfikacjami składniowymi oraz pewnymi ograniczeniami. Restrykcje te dotyczą głównie kodu akcelerowanego, a wśród najważniejszych można wymienić: ograniczenia w alokacji pamięci, brak obsługi wyjątków czy znacznie okrojoną bibliotekę standardową. Program zawierający kod CUDA zwykle umieszcza się w pliku o rozszerzeniu .cu/.cuh w przypadku C/C++ lub .cuf w przypadku Fortran.

Biblioteka CUDA

Część programu CUDA realizowana po stronie hosta może być implementowana za pomocą jednego z dwóch interfejsów:

  • runtime API, wygodniejszy w użyciu interfejs wysokopoziomowy (C/C++ oraz Fortan),
  • driver API, niskopoziomowy interfejs (w języku C).

Obydwa interfejsy mogą być używane jednocześnie. Korzystanie z driver API jest zalecane tylko w wypadku, kiedy konieczne jest użycie funkcjonalności niedostępnych z wyższego poziomu. Przykładowo, za pośrednictwem driver API użytkownik uzyskuje możliwość ręcznego zarządzania tzw. kontekstami CUDA, które w runtime API obsługiwane są automatycznie.

Funkcje z tych dwóch interfejsów można łatwo odróżnić od siebie ─ funkcje z runtime API posiadają przedrostek "cuda", natomiast z driver API "cu". Na poziomie linkowanych bibliotek libcuda odpowiada driver API natomiast libcudart to runtime API.

Model programistyczny

Program CUDA wykonuje się jako zwykły proces, w ramach którego host zleca akceleratorowi różne operacje do zrealizowania. Host jest odpowiedzialny za wykrywanie urządzeń GPU, alokacje i transfery pamięci, inicjowanie wykonania funkcji akcelerowanych (kerneli), a także wszelkie inne funkcjonalności dostępne w ramach biblioteki CUDA. Warto dodać, że wiele różnych procesów może jednocześnie korzystać z tej samej karty GPU.

Delegowanie obliczeń do wykonania na GPU odbywa się poprzez wywoływanie kerneli – funkcji oznaczonych specyfikatorem __global__ (atrybut global w Fortran). W tym celu używa się specjalnej składni kernel_name<<<...>>>(...), która pozwala na określenie stopnia równoległości czyli tego, jak wiele wątków (instancji danego kernela) ma zostać wykonanych równolegle. Wywołanie kernela jest asynchroniczne, co pozwala na kontynuowanie pracy przez hosta i przykładowo prowadzenie swoich obliczeń na CPU albo zlecenie do wykonania kolejnego kernela na tę samą lub inną kartę. Pojedyncza karta umożliwia równoległe wykonanie wielu kerneli (poprzez mechanizm strumieni).

Zarządzanie pamięcią GPU i transferem danych między pamięcią RAM komputera a pamięcią GPU może być realizowane na dwa sposoby:

  • poprzez standardową alokację (cudaMalloc) i jawne wywołania funkcji kopiujących (cudaMemcpy),
  • automatycznie, poprzez zastosowanie mechanizmu Unified Memory w momencie alokacji (cudaMallocManaged).

W pierwszym przypadku programista ma pełną kontrolę nad zachowaniem programu, może również skorzystać z asynchronicznego kopiowania pamięci. W drugim przypadku dane zostają umieszczone we wspólnej, wirtualnej przestrzeni adresowej. W momencie dostępu do nich, automatycznie następuje sprawdzenie czy dane znajdują się fizycznie w pamięci CPU czy GPU i jeśli zachodzi potrzeba, to zostają przerzucone między urządzeniami.

W obu przypadkach należy pamiętać o tym, że transfery pamięci na linii CPU-GPU są operacjami kosztownymi czasowo. W celu wydajnego wykorzystania GPU należy wykonywać je jak najrzadziej lub umiejętnie przeplatać z obliczeniami, tak by wykonywały się równolegle do nich.

Hierarchia wątków

Instancje kernela są zorganizowane w pewną hierarchię ─ wątki są pogrupowane w bloki, te zaś formują grid. Dlatego w składni do uruchomienia kernela podaje się dwa podstawowe argumenty <<<numBlocks, threadsPerBlock>>> oznaczające odpowiednio wymiary gridu (liczbę bloków) oraz wymiary bloku (liczbę wątków w każdym bloku). Argumenty te mogą zostać podane jako pojedyncza liczba lub obiekt przechowujący trzy wartości (dim3), co pozwala na wyrażanie również 2- i 3-wymiarowych gridów. Dodatkowo w najnowszych architekturach została wprowadzona opcjonalna możliwość podziału gridu na klastry zawierające bloki.

Każdy wątek wykonuje ten sam kod, ale posiada swój unikatowy identyfikator w bloku (threadIdx) oraz identyfikator bloku (blockIdx). Pozwala to na rozróżnienie zachowania wątków w kernelu, najczęściej przez prowadzenie tych samych obliczeń ale na danych pobranych z tablicy spod innego indeksu.

poziomy zrownoleglenia a architektura GPU

poziomy zrównoleglenia obliczeń a fizyczna architektura GPU [źródło]

poziomy zrownoleglenia a pamiec

rodzaje pamięci względem hierarchii wątków [źródło]

Podział gridu obliczeniowego na bloki ma praktyczne znaczenie, wynikające z fizycznej architektury GPU – wątki z tego samego bloku są wykonywane przez pojedynczy multiprocesor (sprzętowa jednostka organizacyjna na karcie GPU), a co za tym idzie, mają dostęp do wspólnej pamięci podręcznej i mechanizmów synchronizacji w obrębie swojego bloku. Wątki z różnych bloków prowadzą swoje obliczenia w pełni niezależnie. Ważną cechą obliczeń na GPU jest to, że kolejne wątki wewnątrz bloku są automatycznie grupowane w tzw. warpy o stałej ilości 32 wątków, a multiprocesor (SM ─ streaming multiprocessor) wykonuje naraz tę samą instrukcję dla wszystkich wątków z warpa i zawsze obsługuje je wspólnie.

GPUDirect

W przypadku korzystania z wielu kart GPU w jednym programie (zarówno w obrębie jednego komputera jak i w środowisku rozproszonym) lub intensywnego odczytu danych z dysku, możliwe jest skorzystanie z rozwiązań optymalizujących dostęp do pamięci GPU. Polegają one na możliwości bezpośredniego dostępu do pamięci GPU przez inne urządzenie (inną kartę GPU, dysk, kartę sieciową). Taki dostęp odbywa się bez interakcji z CPU i pamięcią RAM komputera (technika DMA, direct memory access), dzięki czemu transfer danych staje się szybszy.

GPUDirect to rodzina technologii realizujących to założenie. W jej skład wchodzą:

  • GPUDirect P2P (peer to peer),
  • GPUDirect RDMA (remote direct memory access),
  • GPUDirect Storage.

Niektóre z tych rozwiązań są automatycznie dostępne w pakiecie CUDA Toolkit, inne wymagają dodatkowej instalacji.

GPUDirect P2P

Umożliwia bezpośrednie kopiowanie danych lub bezpośredni dostęp (tzw. peer access) przez jeden układ GPU do danych znajdujących się w pamięci innego GPU w obrębie tego samego komputera. Bezpośredni dostęp oznacza, że kernel wykonujący się na jednym GPU może pracować na wskaźnikach do danych z drugiego GPU.

schemat GPUDirect P2P

schemat interakcji GPU i CPU przy kopiowaniu danych między kartami [źródło]

Bezpośredni transfer danych między kartami odbywa się po interfejsie PCI Express (PCIe) lub NVLink i jest możliwy wtedy, jeśli obydwie karty znajdują się w tym samym poddrzewie hierarchii PCIe lub są połączone poprzez NVLink (bądź NVSwitch).

NVLink / NVSwitch

NVLink to fizyczne połączenie i protokół bezpośredniej komunikacji pomiędzy kartami GPU NVIDIA, umożliwiający bardzo szybką wymianę danych (omijającą standardowe łącze PCIe). Wymaga połączenia kart dodatkowym fizycznym mostkiem lub jest realizowane automatycznie w płycie głównej serwera.

NVSwitch to rozszerzenie tej technologii umożliwiające połączenie ze sobą wielu kart GPU w obrębie jednego lub kilku węzłów łączem NVLink (więcej informacji na stronie produktu).

Technologia P2P jest automatycznie dostępna w CUDA Toolkit. Aby z niej skorzystać wystarczy włączyć możliwość bezpośredniego dostępu między parą kart, poprzez wywołanie odpowiedniej funkcji (patrz Peer Device Memory Access).

GPUDirect RDMA

Umożliwia bezpośrednie kopiowanie danych między układami GPU poprzez interfejs sieciowy. Znajduje swoje zastosowanie w przypadku CUDA-aware MPI, gdy dane z pamięci jednej karty GPU są kopiowane do pamięci GPU znajdującego się na innym komputerze.

schemat GPUDirect RDMA

schemat interakcji GPU i CPU przy kopiowaniu danych po sieci [źródło]

Technologia RDMA jest automatycznie dostępna w CUDA Toolkit. Jej wykorzystanie jest zależne od tego czy karta sieciowa i jej sterownik wspierają tę technologię. Po stronie programistycznej skorzystanie z niej jest automatycznie i nie wymaga żadnych dodatkowych funkcji.

GPUDirect Storage

Umożliwia bezpośrednie wczytywanie i zapisywanie danych z dysku (zarówno lokalnego jak i sieciowego). Technologia GPUDirect Storage nie jest dostępna w standardowej wersji. Wymaga posiadania dysków i sterowników przystosowanych do tego trybu pracy oraz instalacji dodatkowego pakietu.

schemat GPUDirect Storage

schemat interakcji GPU i CPU przy pobieraniu danych z dysku [źródło]

Terminologia

  • host ─ Główny proces programu CUDA, wykonywany na CPU, delegujący obliczenia do wykonania na akceleratorach.

  • akcelerator (device) ─ Układ GPU.

  • kernel ─ Funkcja przeznaczona do wykonania na akceleratorze, wywoływana przy użyciu specjalnej składni <<<...>>>.

  • konfiguracja wykonania ─ Parametry wywołania kernela, podawane w składni <<<...>>>. Podstawowo służy do określenia liczby bloków (wymiarów gridu) oraz liczby wątków w bloku (wymiarów bloku). Dodatkowo (opcjonalnie) pozwala określić rozmiar dynamicznie alokowanej pamięci współdzielonej dla bloku, oraz wskazać strumień (inny niż domyślny) w którym ma zostać wykonany kernel.

  • strumień (stream) ─ Mechanizm asynchronicznego równoległego wykonywania kilku kerneli lub transferów pamięci na tym samym GPU. Operacje zlecone do różnych strumieni mogą zostać zrealizowane w dowolnej kolejności, w szczególności mogą wykonywać się równolegle. Natomiast pojedynczy strumień funkcjonuje jak kolejka FIFO, tzn. operacje zlecone do niego będą uruchamiane w kolejności w jakiej zostały wywołane.

  • wątek, blok, klaster, grid (thread, thread block, thread block cluster, grid) ─ Poziomy organizacji obliczeń na układach GPU w CUDA:

    • wątek ─ reprezentuje pojedynczą instancję kernela,
    • blok ─ grupa wątków wykonywanych na tym samym multiprocesorze, posiadająca dostęp do wspólnej, szybkiej pamięci (shared memory), mogąca być synchronizowana,
    • klaster ─ dodatkowy poziom organizacji obliczeń dostępny w najnowszych układach GPU (od compute capability 9.0), pozwalający zgrupować ze sobą kilka bloków, dać im dostęp do wspólnej pamięci (distributed shared memory) oraz możliwość synchronizacji (bloki z klastra będą wykonywane w tej samej jednostce GPC, zawierającej kilka multiprocesorów),
    • grid ─ grupa wszystkich bloków luk klastrów wykonujących dany kernel, realizowana przez cały układ GPU (różne bloki mogą być wykonywane na różnych multiprocesorach).

  • warp ─ Grupa wątków z tego samego bloku przetwarzanych jednocześnie przez multiprocesor (architektura SIMT), tj. wspólnie zarządzanych i mogących w danym momencie wykonywać równolegle tylko jedną wspólną instrukcję (każdy wątek może podążać swoją ścieżką obliczeń, multiprocesor decyduje którą operację w danym momencie wykonać i zostaje ona zrealizowana tylko przez te wątki, które aktualnie czekają właśnie na tę instrukcję). W architekturze CUDA warp zawsze liczy 32 wątki. Podział na warpy odbywa się automatycznie, wedle kolejności wątków w bloku. Istnieją specjalne funkcje dla operacji wewnątrz warpa.

  • funkcje __host__, __global__, __device__- Rodzaje funkcji, określane poprzez umieszczenie odpowiedniego specyfikatora przy deklaracji funkcji (w języku Fortran poprzez atrybuty: host, global, device):

    • host ─ standardowe funkcje, wykonywane przez CPU,
    • global ─ kernele wykonywane na GPU, możliwe do wywołania z kodu CPU, a także z innego kernela (mechanizm CUDA Dynamic Parallelism),
    • device ─ funkcje pomocnicze wykonywane na GPU, możliwe do wywołania z kernela lub innej funkcji device.

    Domyślnie każda funkcja jest typu host. Możliwe jest oznaczenie funkcji jednocześnie jako host oraz device, co pozwala na wywoływanie jej zarówno z kodu zwykłego jak i kodu akcelerowanego.

  • zmienne gridDim, blockDim ─ Specjalne zmienne dostępne w kernelu, pozwalające uzyskać informacje o wymiarach gridu obliczeniowego dla aktualnego wywołania kernela. Odpowiadają wartościom przekazywanym w konfiguracji wykonania kernela.

  • zmienne threadIdx, blockIdx ─ Specjalne zmienne dostępne w kernelu, pozwalające uzyskać informacje o położeniu wątku w bloku, oraz bloku w całym gridzie. Na ich podstawie wątek może zróżnicować swoje zachowanie i prowadzić obliczenia na własnej porcji danych.

  • register/local/shared/global/constant memory ─ Podstawowe rodzaje pamięci dostępne na GPU. Możliwe do określenia przez dodanie odpowiednich specyfikatorów.

    • register ─ Pamięć podręczna przeznaczona dla każdego z wątków z osobna. Cechuje ją najmniejsza objętość oraz największa przepustowość z najmniejszym opóźnieniem. W pamięci register automatycznie umieszczane są zmienne lokalne kernela.

    • local ─ Pamięć przeznaczona dla każdego z wątków z osobna. Automatycznie są w niej umieszczane zmienne lokalne kernela, w sytuacji gdy nie mieszczą się one w pamięci typu register (przesłanki kiedy tak się dzieje ─ patrz Device Memory Accesses). Jest ona wydzielona z pamięci global, a więc jest dużo wolniejsza niż pamięć register.

    • shared ─ Pamięć podręczna dla danego kernela, współdzielona przez wszystkie wątki z tego samego bloku, znacząco szybsza od pamięci global. Jej użycie zależy bezpośrednio od programisty ─ aby tablica z danymi została umieszczona w pamięci shared należy zadeklarować ją z użyciem specyfikatora __shared__ (atrybut shared w Fortran). Dodatkowo, jeśli kilka bloków zostaje zgrupowanych w klaster, zyskują one dostęp do tzw. distributed shared memory, czyli połączonej pamięci typu shared wszystkich bloków w danym klastrze.

    • global ─ Główna pamięć układu GPU, odpowiednik pamięci RAM dla CPU. Posiada dużo większą pojemność (gigabajty) ale dużo mniejszą wydajność niż pamięć register czy shared. Najczęściej alokowana z poziomu hosta i dostępna do użytku z poziomu wszystkich wątków wykonujących dany kernel. W runtime API i kernelach dla C/C++ adresowana przez zwykłe wskaźniki. Transfer danych między pamięcią global a pamięcią hosta jest możliwy poprzez funkcje kopiujące (cudaMemcpy) lub automatycznie w przypadku mechanizmu Unified Memory. Dodatkowo, poszczególne wątki oraz bloki wątków mogą alokować własną przestrzeń w pamięci global (patrz Dynamic Global Memory).

    • constant ─ Niewielki fragment pamięci (ograniczony do 64KB), który jest dostępny w trybie tylko do odczytu przez wszystkie wątki wykonujące dany kernel. Kolejne odczyty trafiają do pamięci podręcznej, dzięki czemu jest to najbardziej efektywne miejsce do umieszczenia stałych danych, używanych przez wiele wątków.

Informacje o standardzie

Model CUDA jest dojrzały i bardzo rozbudowany. Jego początki sięgają 2007 roku i od tego momentu jest stale rozwijany przez firmę NVIDIA, równolegle do rozwoju własnych architektur sprzętowych. Stanowi podstawę szerokiego ekosystemu programistycznego i jest wydawany w formie pakietu CUDA Toolkit lub rozszerzonego NVIDIA HPC SDK. W ich skład wchodzą biblioteki gotowe do pracy w modelu CUDA, inne modele programowania GPU (np. OpenACC), oraz rozbudowane narzędzia do analizy wydajności i poprawności kodu.

CUDA była pierwszym rozwiązaniem przeznaczonym do wykorzystania układów GPU do obliczeń ogólnego przeznaczenia GPGPU (general-purpose computing on graphics processing units). Dzięki wieloletniej dominacji firmy NVIDIA na rynku kart graficznych, CUDA stała się najpopularniejszym modelem programowania GPU i de facto wyznaczyła standard programowania w tym obszarze. Późniejsze rozwiązania (OpenCL, SYCL, HIP) powstawały w odniesieniu do modelu CUDA i zawierają wiele analogii. Obecnie, dzięki rozwojowi kart GPU firmy AMD oraz Intel, inne technologie programowania tego typu urządzeń zaczynają zyskiwać na znaczeniu.

CUDA Toolkit

Podstawowy pakiet oprogramowania i narzędzi potrzebnych do tworzenia, uruchamiania i analizowania aplikacji korzystających z CUDA. Zawiera m.in.:

  • kompilator nvcc,
  • biblioteki uruchomieniowe CUDA,
  • podstawowe narzędzia (profiler, debugger),
  • biblioteki programistyczne wspierające CUDA.

Aktualna wersja to CUDA Toolkit 12.0.1, wydana w styczniu 2023. Spis zawartości pakietu wraz z numerami wersji poszczególnych komponentów znajduje się w informacjach o wydaniu. Spośród najważniejszych zmian względem wersji 11.x można wymienić dodanie wsparcia dla nowych architektur GPU (Ada Lovelace oraz Hopper), a także wprowadzenie możliwości korzystania z niektórych funkcjonalności języka C++20.

Numeracja pakietów odpowiada rozwojowi możliwości programistycznych i wiąże się z zagadnieniem kompatybilności:

  • wersje major są wydawane zwykle przy okazji premiery nowej architektury procesorów GPU i często usuwają wsparcie dla starszych układów (np. w wersji 12.0 zakończono wsparcie dla architektury Kepler);
  • wersje minor (np. 11.1, 11.2) dodają nowe funkcjonalności lub wsparcie dodatkowych układów, zachowując przy tym kompatybilność w obrębie wersji major.

Aktualizacje pakietu ukazują się kilka razy w roku. Obecna wersja CUDA Toolkit wspiera urządzenia o compute capability od 5.0 do 9.0 (co odpowiada architekturom Maxwell, Pascal, Volta, Turing, Ampere, Ada Lovelace, Hopper).

Compute Capability

Zakres wspieranych architektur sprzętowych przez daną wersję CUDA Toolkit jest określany przez wartość compute capability. Jest to oznaczenie podstawowych zdolności obliczeniowych (specyfikacji technicznej) oraz funkcjonalności odpowiednich dla kolejnych generacji kart GPU firmy NVIDIA (patrz wykaz kart GPU wedle compute capability).

Duże zmiany w compute capability pojawiają się wraz z kolejnymi architekturami układów GPU (które noszą różne nazwy, np. Volta, Ampere, Hopper), natomiast w obrębie jednej architektury mogą występować mniejsze różnice (np. compute capability 8.0, 8.6, 8.7 odnoszą się do różnych układów o architekturze Ampere).

NVIDIA HPC SDK

Rozszerzony pakiet narzędzi programistycznych i bibliotek, dostosowany do pracy w środowisku HPC, z uwzględnieniem wykorzystania wielu kart GPU w modelu rozproszonym. W jego skład, poza elementami pakietu CUDA Toolkit, wchodzą:

  • dodatkowe kompilatory (nvc, nvc++, nvfortran), umożliwiające
    • wykorzystanie modelu CUDA w języku Fortran,
    • korzystanie z innych modeli programowania GPU (OpenACC, OpenMP offloading, C++ Parallel Algorithms) zarówno w języku C++ jak i Fortran,
  • rozwiązania dotyczące pracy w środowisku rozproszonym (OpenMPI, HPC-X, UCX, NVSHMEM) i wymiany danych między wieloma kartami (NCCL),
  • dodatkowe biblioteki (cuTENSOR) oraz wersje podstawowych bibliotek przystosowane do pracy w architekturze rozproszonej (cuFFTMp, cuSOLVERMp),
  • dystrybucje podstawowych bibliotek obliczeniowych dla CPU (OpenBLAS, Scalapack).

Wydawany kilka razy w roku niezależnie od podstawowego pakietu. Aktualna wersja to HPC SDK 23.1, wydana w styczniu 2023. Bazuje na CUDA Toolkit 12.0, dostępny jest też wariant zawierający trzy poprzednie wersje (12.0, 11.8 oraz 11.0). Spis zawartości pakietu wraz z numerami wersji poszczególnych komponentów znajduje się w informacjach o wydaniu.

Sterownik

Aby uruchamiać oprogramowanie korzystające z CUDA konieczne jest posiadanie zainstalowanego sterownika NVIDIA (NVIDIA display driver) w odpowiedniej wersji. Każda wersja CUDA Toolkit określa minimalną wymaganą wersję NVIDIA driver. Aktualna wersja sterownika dostarczana jest jako element CUDA Toolkit (pakiet HPC SDK nie zawiera jej), ale może być też zainstalowana oddzielnie (na przykład za pomocą menedżera pakietów w dystrybucjach systemu Linux lub poprzez pobranie ze strony).

schemat srodowiska CUDA

składowe środowiska CUDA [źródło]

CUDA umożliwia uruchamianie aplikacji zbudowanych w nowszym środowisku na systemie posiadającym starszy sterownik. Informacje na ten temat można znaleźć w dokumentacji dotyczącej kompatybilności.

Kompilatory

Aby skompilować program CUDA potrzebny jest specjalny kompilator, który będzie potrafił zinterpretować składnię bazowego języka rozszerzoną o elementy CUDA, oraz będzie potrafił wygenerować kod dla GPU. Służą do tego oficjalne kompilatory dostarczane przez firmę NVIDIA (nvcc, nvfortran), jest również dostępna alternatywa w postaci kompilatora clang.

Rolą kompilatora jest rozbicie kodu CUDA na kod hosta i kod akcelerowany, kompilacja każdego z nich i połączenie w jeden plik wykonywalny lub bibliotekę. Kod akcelerowany jest kompilowany pod wskazane compute capability, ten sam kod może zostać równocześnie skompilowany pod wiele architektur.

Szczegóły postaci binarnej

Efektem końcowym kompilacji kodu akcelerowanego może być postać binarna (cubin) lub kod pośredni (w formacie PTX), zgodne z danym compute capability. Następnie postać cubin/ptx zostaje umieszczona w tzw. fat binary (fatbin), który umożliwia przechowywanie wielu wersji cubin/ptx dla tego samego kodu. Domyślnie kod cubin/ptx jest umieszczany w postaci fatbin wraz z kodem binarnym dla CPU wewnątrz jednego standardowego pliku wykonywalnego lub biblioteki. Można jednak poinstruować kompilator aby kod akcelerowany został zapisany w formie osobnych plików cubin/ptx/fatbin. Takie pliki mogą zostać wykorzystane w innym programie CUDA za pomocą driver API (patrz Module Management np. funkcja cuModuleLoad).

W przypadku umieszczenia w pliku wykonywalnym kilku wersji ptx/cubin tego samego kodu, runtime CUDA w trakcie wykonania programu automatycznie wybiera wersję najbardziej odpowiednią dla używanego GPU (reguły wyboru ─ patrz kompatybilność aplikacji). Dzięki różnym wersjom skompilowanego kodu akcelerowanego możliwe jest pełne wykorzystanie funkcjonalności kilku różnych architektur GPU oraz zapewnienie możliwości wykonania na nowszych urządzeniach. Odbywa się to za cenę większego rozmiaru końcowego pliku binarnego.

cubin

Kod skompilowany do postaci binarnej cubin jest natywną, najniżej poziomową postacią programu akcelerowanego i można go uruchomić tylko na urządzeniach o konkretnym, określonym podczas kompilacji, compute capability. Taki kod binarny jest kompatybilny tylko z nowszymi urządzeniami w ramach tej samej architektury (tj. z urządzeniami o tym samym głównym numerze compute capability, ale o większym drugim numerze).

Przykładowo dla kodu binarnego w wersji 8.6 możliwe jest jego uruchomienie na urządzeniach o compute capability 8.7 i 8.9, ale na urządzeniach z compute capability 9.0 już nie można go uruchomić. Więcej informacji ─ patrz kompatybilność binarna. W niektórych opisach tego zagadnienia można natknąć się na określenie SASS, które oznacza język asemblerowy dla postaci binarnej.

PTX

Kod pośredni PTX (Parallel Thread Execution) jest niskopoziomowym zapisem instrukcji programu akcelerowanego dla wirtualnej maszyny oddającej zachowanie GPU. Jak nazwa wskazuje, nie jest kodem gotowym do bezpośredniego wykonania na GPU ─ niezbędna jest jego kompilacja do kodu binarnego. Składnia PTX jest publicznie udokumentowana.

Podczas działania programu z częścią akcelerowaną w postaci PTX, kod pośredni jest kompilowany w locie (just-in-time compilation) do wersji binarnej odpowiedniej dla używanego GPU. Domyślnie dzieje się to w momencie uruchomienia programu ale można zmienić to zachowanie aby kod był kompilowany tylko dla używanych kerneli w momencie ich pierwszego użycia (tzw. lazy loading). Powoduje to pewien narzut czasowy, ale skompilowana wersja binarna w formacie cubin zostaje umieszczona w pamięci, dzięki czemu ten dodatkowy koszt czasowy jest zwykle jednorazowy. Wszystkie te operacje są wykonywane automatycznie przez runtime CUDA.

Takie podejście służy zwiększeniu kompatybilności i umożliwia uruchamianie raz skompilowanego kodu również na urządzeniach o wyższym compute capability niż określone podczas pierwotnej kompilacji, oraz umożliwia wykorzystanie optymalizacji wprowadzonych w nowszych kompilatorach (należy jednak pamiętać, że kod PTX zgodny z wcześniejszym compute capability może nie być w stanie wykorzystać wszystkich funkcjonalności wprowadzonych w nowszych compute capability). Więcej informacji ─ patrz kompatybilność PTX.

Opis kompilatorów

Poniżej znajdują się opisy kompilatorów:

NVCC

Podstawowy kompilator kodu CUDA dla języków C oraz C++. Do prawidłowego działania wymaga dostępu do jednego z wspieranych kompilatorów C/C++ (gcc, icc, nvc/nvc++, clang, msvc). W trackie standardowej kompilacji, dzieli kod źródłowy na część akcelerowaną oraz nieakcelerowaną. Kompilacja kodu akcelerowanego wykonywana jest bezpośrednio przez nvcc, natomiast kod hosta najpierw zostaje przetłumaczony na kod C/C++ (bez rozszerzeń składniowych CUDA), a następnie zostaje skompilowany standardowym kompilatorem C/C++.

Domyślnie nvcc przeprowadza kompilację do końcowego pliku wykonywalnego lub biblioteki. Umożliwia wykonanie tylko wskazanych faz kompilacji ─ przykładowo zapis kodu akcelerowanego do samodzielnych plików .cubin, .ptx lub .fatbin, czy kompilacja z tych plików do dalszej postaci. Pracując z różnymi plikami wejściowymi interpretuje je na podstawie rozszerzenia.

Podstawowe opcje kompilacji umożliwiają określenie dla jakiej architektury ma zostać skompilowany kod, oraz czy ma zostać wygenerowana postać cubin czy PTX. Rozróżniane są dwa rodzaje architektur (numerowane według compute-capability): rzeczywiste (sm_) oraz wirtualne (compute_), odpowiadające postaci natywnej (cubin) i pośredniej (PTX). Odpowiednie opcje kompilacji to:

  • -arch, wskazuje wirtualną architekturę, podana samodzielnie definiuje również w jakiej postaci ma zostać wygenerowany kod;
  • -code, odpowiada za specyfikację postaci binarnej (PTX lub cubin), przy czym podana architektura musi być kompatybilna z architekturą podaną w -arch.
Przykłady opcji kompilacji
  • -arch=compute_80 spowoduje wygenerowanie PTX zgodnego z compute capability 8.0
  • -arch=sm_80 spowoduje wygenerowanie PTX oraz cubin zgodnego z compute capability 8.0
  • -arch=compute_86 -code=sm_86, wygeneruje tylko cubin zgodny z compute capability 8.6

Do tworzenia bardziej wyspecjalizowanych kombinacji (np. do generowania kodu pod różne architektury wirtualne) można wykorzystać flagę -gencode.

Kompilator nvcc wchodzi w skład podstawowego pakietu CUDA Toolkit, jego obecna wersja to 12.0.140, wydana w styczniu 2023 w ramach CUDA Toolkit 12.0.1. Jest również dostępny w kilku wersjach w ramach rozszerzonego pakietu NVIDIA HPC SDK.

Więcej informacji

NVC/NVC++

Kompilatory języków C i C++, które oprócz klasycznej kompilacji kodu na CPU, umożliwiają wykorzystanie kart GPU NVIDIA przy użyciu innych modeli niż CUDA C/C++, takich jak: OpenMP, OpenACC oraz C++17 Parallel Algorithms (akceleracja algorytmów biblioteki standardowej na układach GPU). Nie wspierają klasycznego modelu CUDA C/C++, tj. nie potrafią kompilować ręcznie napisanych kerneli, natomiast w pełni nadają się na standardowy kompilator uzupełniający pracę nvcc.

Wraz z nvfortran wchodzą w skład NVIDIA HPC SDK i są nazywane kompilatorami HPC ─ są zbliżone do siebie i w dużej mierze wspierają te same opcje kompilacji. Obecna wersja to 23.1 (zgodna z numerem pakietu HPC SDK), wydana w styczniu 2023. Kompilatory NVC/NVC++ dawniej były rozwijane pod szyldem PGI (pgcc/pgc++).

Więcej informacji

NVFORTRAN

Podstawowy kompilator kodu CUDA dla języka Fortran, a jednocześnie w pełni samodzielny kompilator języka Fortran (w wersji 2003 oraz częściowo 2008). Umożliwia również wykorzystanie kart GPU NVIDIA przy użyciu innych technologii niż model CUDA, takich jak: OpenMP, OpenACC oraz standardowe elementy języka Fortran wyrażające równoległość obliczeń (konstrukcja do concurrent). Jest to więc odpowiednik NVCC połączonego z NVC++ dla języka Fortran.

Możliwości CUDA dla języka Fortran odpowiadają możliwością dla języka C/C++, chociaż należy mieć na uwadze, że niektóre funkcjonalności mogą być dostępne tylko w bazowym modelu. Z drugiej strony nvfortran posiada wsparcie dla dodatkowych dyrektyw !@cuf, pozwalających na oznaczenie instrukcji które mają być kompilowane tylko w przypadku korzystania z CUDA (kompilacja warunkowa) oraz !$cuf kernel pozwalających na automatyczne generowanie kerneli ze standardowych pętli (funkcjonalność Kernel Loop Directive).

Obsługa nvfortan różni się nieco od nvcc. Do kompilowania kodu akcelerowanego konieczne jest przekazanie flagi -cuda lub umieszczenie kodu w pliku z rozszerzeniem .cuf (rozszerzenia plików wejściowych wpływają na zachowanie kompilatora). Z kolei flaga -gpu pozwala określić szczegóły generacji kodu, przykładowo:

  • -gpu=ccXY pozwala wyspecyfikować docelowe compute capability (o numerze X.Y),
  • -gpu=cudaX.Y pozwala określić kompatybilność ze wskazanym CUDA Toolkit.

Domyślnie nvfortran włącza do końcowego pliku zarówno postać binarną skompilowanego kodu akcelerowanego jak i kod pośredni PTX (kod PTX nie jest dołączany w przypadku kompilacji z flagą -⁠gpu=nordc). Kompilator nie pracuje bezpośrednio z plikami ptx/cubin/fatbin.

Kompilator nvfortran wraz z nvc/nvc++ wchodzą w skład NVIDIA HPC SDK i są nazywane kompilatorami HPC ─ są zbliżone do siebie i w dużej mierze wspierają te same opcje kompilacji. Obecna wersja to 23.1 (zgodna z numerem pakietu HPC SDK), wydana w styczniu 2023. Wspiera compute capability od 3.5 do 8.6. Dawniej był rozwijany pod szyldem PGI (pgfortran).

Więcej informacji

Clang

Otwartoźródłowy kompilator języków C/C++ potrafiący kompilować kod CUDA C/C++. Jego rozwój nie jest tak szybki jak flagowego kompilatora, ale stanowi realną alternatywę dla nvcc. Do korzystania z cuda-clang wciąż wymagany jest dostęp do CUDA Toolkit ─ w celu linkowania z bazowymi bibliotekami (libcuda, libcudart). Deweloperzy clang nie gwarantują działania kompilacji kodu CUDA na systemach operacyjnych innych niż Linux.

W odróżnieniu od nvcc, clang samodzielnie wykonuje całość kompilacji, zarówno kodu hosta jak i kodu akcelerowanego. Aby kompilacja CUDA została włączona, kod musi być umieszczony w pliku o rozszerzeniu .cu lub trzeba podać opcję -x cuda. Wynikiem kompilacji jest końcowy plik wykonywalny zawierający część akcelerowaną w postaci PTX. Podobnie jak w przypadku nvcc, możliwe jest kompilowanie pod kątem wskazanych compute capability, przy czym służy do tego inna flaga kompilacji (--cuda-gpu-arch=sm_XX). W zależności od rodzaju aplikacji kod wyprodukowany przez clang może być porównywalnej (bądź lepszej) wydajności a sam proces kompilacji może wykonać się szybciej w stosunku do kompilatora nvcc.

Clang jest rozwijany w ramach projektu LLVM. Obecna wersja to 15.0.7, wydana w styczniu 2023. Jest kompatybilna z CUDA do wersji 11.5, tj. nie potrafi korzystać z funkcjonalności specyficznych dla nowszych wersji, ale w większości przypadków powinna działać również z nowszymi wydaniami (traktując je jak 11.5). Nowe główne wersje Clang pojawiają się co ok. pół roku, wsparcie dla CUDA jest aktywnie rozwijane.

Więcej informacji

Narzędzia

Nvidia udostępnia szeroką gamę narzędzi usprawniających programistom pracę z aplikacjami CUDA. Można je odnaleźć pod nazwą “NVIDIA Nsight tools”. Poniżej zostały wymienione podstawowe spośród nich. Więcej informacji oraz pełna lista tych narzędzi ─ patrz NVIDIA Developer Tools.

  • NVIDIA Nsight Systems ─ podstawowy profiler dla oprogramowania korzystającego z CUDA. Umożliwia całościową analizę wydajności aplikacji ukazując ile czasu jest spędzane na obliczeniach CPU, ile na GPU, ile na transferach pamięci. Profil działania aplikacji (timeline) pozwala zobaczyć kiedy jaka funkcja/kernel były wykonane oraz jaka była ich wydajność. Dla obliczeń na GPU rejestruje różne metryki, które pozwalają na podstawową analizę kerneli pod kątem wydajności.
    https://developer.nvidia.com/nsight-systems

  • NVIDIA Nsight Compute ─ zaawansowany profiler CUDA. Potrafi dostarczyć szczegółowych informacji o wykorzystaniu zasobów sprzętowych GPU, pamięci, zachowaniu API CUDA i wiele innych. Pozwala na dogłębną analizę wykonania kerneli, umożliwiając optymalizację pod kątem możliwości sprzętowych posiadanego układu GPU.
    https://developer.nvidia.com/nsight-compute

  • cuda-gdb ─ debugger bazujący na gdb (standardowy debugger GNU) umożliwiający pracę z aplikacjami CUDA. Pozwala na wykrywanie błędów, przerywanie działania programu, wykonywanie kodu krok po kroku, podglądanie wartości zmiennych, stosu wywołań funkcji itp.
    https://developer.nvidia.com/cuda-gdb

Oprócz tego są dostępne również inne narzędzia, przy czym niektóre z nich nie są już aktualne.

  • CUDA binary utils ─ zbiór narzędzi (cuobjdump, nvdisasm, cu++filt, nvprune) do pracy z plikami binarnymi zawierającymi kod akcelerowany w postaci PTX/cubin.
    https://docs.nvidia.com/cuda/cuda-binary-utilities/

  • ⚠️ CUDA Occupancy Calculator ─ kalkulator wykorzystania zasobów GPU w formie arkusza kalkulacyjnego. Obecnie nie jest aktualizowany, a analogiczna funkcjonalność została zaimplementowana w oprogramowaniu Nsight Compute. Może być jednak wciąż wykorzystywany dla GPU do compute capability 8.6.
    https://docs.nvidia.com/cuda/cuda-occupancy-calculator/index.html

  • ⚠️ NVIDIA Visual Profiler / nvprof ─ pierwotny profiler dla aplikacji CUDA. Obecnie zastąpiony przez narzędzia NVIDIA Nsight. Proces migracji jest opisany w artykule oraz dokumentacji.

Linki

  1. Pierwotnie nazwa CUDA pochodzi od Compute Unified Device Architecture. Obecnie już nie jest stosowana jako akronim. ↩

Ostatnia aktualizacja: 15 września 2023