Przejdź do treści

HIP

logo ROCm

Wprowadzenie

HIP (Heterogeneous-Computing Interface for Portability) to rozwijany przez firmę AMD model programowania dla platform akcelerowanych będący odpowiednikiem NVIDIA CUDA. Umożliwia tworzenie oprogramowania dla układów GPU AMD, ale jednocześnie gwarantuje przenośność kodu, tj. pozwala na uruchomienie tego samego kodu również na GPU NVIDIA. Wchodzi w skład szerszego zestawu narzędzi programistycznych oraz bibliotek AMD ROCm.

Interfejs programistyczny HIP ściśle bazuje na CUDA. Można powiedzieć, że jest to podzbiór CUDA z pewnymi modyfikacjami wprowadzonymi przez AMD. Z perspektywy programisty różnice w znacznej mierze sprowadzają się tylko do zmiany prefiksu nazw funkcji z "cuda..." na "hip...". Takie rozwiązanie pozwala na stosunkowo szybkie i łatwe przenoszenie aplikacji z CUDA do HIP. Dla ułatwienia tego procesu AMD dostarcza narzędzie HIPify. Należy jednak pamiętać, że pewne funkcjonalności CUDA nie są (obecnie) wspierane przez HIP i ich przeniesienie może być niemożliwe lub wymagać większego nakładu pracy.

Dostępność

Podstawowym językiem programowania dla modelu HIP jest C++. Podobnie jak CUDA, HIP oparty jest o rozszerzenie składni języka oraz funkcje biblioteczne. Wymaga specjalnego kompilatora (hipcc lub amdclang++) oraz dostępu do zainstalowanego pakietu AMD ROCm (więcej w sekcji informacje o standardzie). Podobnie, w przypadku kompilacji pod kątem GPU NVIDIA, HIP wymaga także zainstalowanego środowiska CUDA (więcej w przenośność kodu).

HIP może również być używany z językiem Fortran, ale w ograniczony sposób. Służy do tego moduł hipfort, który umożliwia korzystanie z API HIP oraz bibliotek HIP/ROCm (takimi jak np. hipBLAS/rocBLAS, hipFFT/rocFFT).

hipfort a CUDA Fortran

Moduł hipfort nie jest odpowiednikiem CUDA Fortran, w szczególności nie umożliwia pisania kerneli bezpośrednio w języku Fortran. Kernele należy pisać w C++, opakowywać w funkcje C, po czym wywoływać je z kodu Fortran przez standardowy mechanizm łączenia tych dwóch języków. Więcej informacji w prezentacji HIPFort: Present and Future Directions for Portable GPU Programming in Fortran (dostępne jest nagranie).

Inne technologie

HIP może być używany razem z innymi technologiami pozwalającymi na zrównoleglenie lub rozproszenie obliczeń takimi jak OpenMP czy MPI. Dodatkowo kompilatory HIP z pakietu ROCm wspierają wykorzystanie kart GPU przy użyciu OpenMP offloading oraz OpenCL. Z kolei rozwiązanie GPU-Enabled MPI (odpowiednik CUDA-aware MPI, wykorzystywana jest również nazwa ROCm-aware MPI) umożliwia efektywne korzystanie z MPI, poprzez przesyłanie danych bezpośrednio z pamięci GPU AMD (technologia RDMA, tj. remote direct memory access).

Istnieje również odpowiednik rozwiązania NVIDIA NCCL – biblioteka RCCL (ROCm Communication Collectives Library). Udostępnia ona zbliżone do MPI interfejsy komunikacji oraz wymiany danych pomiędzy wieloma GPU w modelu rozproszonym. RCCL może być wykorzystywany zarówno z kartami GPU w ramach jednego systemu jak i wieloma węzłami obliczeniowymi wyposażonymi w akceleratory GPU wspierane przez HIP/ROCm.

Szczegóły

Model programistyczny HIP, a także zbiór bibliotek wchodzących w skład pakietu ROCm, w przeważającej większości są wiernymi odwzorowaniami rozwiązań znanych z CUDA. Z tego powodu poniższy tekst skupia się głównie na różnicach między tymi modelami. Bazowe założenia są zarysowane tylko ogólnie, zakładając podstawową znajomość CUDA (patrz CUDA > Szczegóły).

HIP wykorzystuje klasyczną strukturę kodu akcelerowanego z kodem hosta (CPU) oraz kernelami, czyli procedurami przeznaczonymi do wykonania na akceleratorze (GPU). Host deleguje kernele do GPU i kontynuuje swoją pracę, co umożliwia równoczesne prowadzenie obliczeń zarówno na CPU jak i GPU. Program zawierający kod HIP zwykle umieszcza się w pliku o rozszerzeniu .hip albo .hip.cpp / .hip.h (według rekomendacji AMD dot. nazwy pliku).

Interfejs HIP

Interfejs HIP zasadniczo odpowiada runtime API z modelu CUDA. Poszczególne funkcje odpowiadają sobie, przykładowo hipMalloc to odpowiednik cudaMalloc. Funkcjonalność HIP jest podzbiorem funkcjonalności CUDA. Warto nadmienić, że nie wszystkie rozwiązania z CUDA są dostępne z poziomu HIP. Spis większości dostępnych funkcji można znaleźć sprawdzając porównanie API CUDA względem HIP.

HIP zawiera również niektóre funkcje odpowiadające driver API z CUDA. Pozwalają one na wyrażenie niżej poziomowych operacji dotyczących GPU.

Model programistyczny

HIP stosuje podejście i składnię znaną z CUDA. Kernele są deklarowane przez dodanie specyfikatora __global__ przy definicji funkcji. Można je wywoływać na dwa sposoby:

  • korzystając ze standardowej dla CUDA składni kernel_name<<<numBlocks, threadsPerBlock>>>(...),
  • poprzez polecane przez AMD hipLaunchKernelGGL.

Funkcja hipLaunchKernelGGL to rozwiązanie niewykraczające poza składnię języka C/C++. Przyjmuje jako argumenty: nazwę kernela, konfigurację wykonania (liczba bloków, liczba wątków w bloku, rozmiar pamięci dynamicznej, identyfikator strumienia), oraz argumenty kernela. Dostępna jest również funkcja hipLaunchKernel o składni zgodnej z C.

Przykład wywołania kernela 
__global__ void MyKernel(float *A, float *B, float *C, size_t N) { ... }

// wariant 1:
MyKernel<<<dim3(gridDim), dim3(groupDim), 0, 0>>> (a, b, c, n);

// wariant 2:
hipLaunchKernelGGL(MyKernel, dim3(gridDim), dim3(groupDim), 0, 0, a, b, c, n)
Przestarzałe konwencje wołania kerneli

W niektórych przykładach dotyczących HIP można znaleźć wzmiankę o parametrze hipLaunchParm (wymaganym jako pierwszy argument kernela). Wydaje się, że jest to przestarzały sposób deklarowania kerneli. Obecnie nie jest potrzebny, a kod w starej składni może się nie kompilować.

Zarządzanie pamięcią GPU odbywa się na dwa sposoby:

  • ręcznie, przez hipMalloc oraz jawne kopiowania hipMemcpy,
  • automatyczne, poprzez zastosowanie hipMallocManaged.

Wywołanie kernela odpowiada wykonaniu wielu instancji danej funkcji w formie 3-wymiarowego gridu obliczeniowego. Grid obliczeniowy jest podzielony na bloki, każdy grupuje wiele wątków (pojedynczych instancji). Tak samo jak w CUDA, w kodzie funkcji akcelerowanych stosuje się zmienne threadIdx oraz blockIdx, które jednoznacznie identyfikują dany wątek.

warp (wavefront) na platformach AMD oraz NVIDIA

Na poziomie wykonywania obliczeń na GPU kolejne wątki są automatycznie łączone w grupy, które wykonują obliczenia jednocześnie – znane z CUDA jako warpy. To samo pojęcie jest wykorzystywane obecnie w HIP, ale w starszych wersjach nosiło nazwę wavefront.

Jedną z istotnych różnic między architekturami jest to, że w przypadku GPU AMD warp liczy 64 wątki. Tymczasem w CUDA (czyli na GPU NVIDIA) warp liczy 32 wątki. Oznacza to, że – w zależności od tego na jakim sprzęcie będzie uruchomiony – ten sam kod niskopoziomowo będzie zachowywał się inaczej lub nie będzie optymalnie wykorzystywał zasobów sprzętowych. Należy wziąć to pod uwagę optymalizując aplikację pod konkretną architekturę bądź korzystając z funkcji wykonywanych przez wszystkie wątki z warpa (tzw. warp cross-lane functions).

Przenośność kodu

W przeciwieństwie do modelu CUDA (wspierającego tylko urządzenia NVIDIA), HIP ma gwarantować przenośność kodu między architekturami – kod powstały z użyciem HIP może być kompilowany i uruchamiany zarówno na kartach GPU AMD jak i NVIDIA. Jest to możliwe dzięki kompilatorowi hipcc . Udostępnia ono dwie ścieżki kompilacji: dla platform AMD, oraz dla platform NVIDIA.

Rozwój standardu

Schemat kompilacji HIP na różnych platformach GPU [źródło]

W przypadku kompilacji pod kątem GPU AMD, hipcc wywołuje wewnętrznie kompilator amdclang++ (potocznie nazywany również HIP-Clang), który odpowiada za faktyczny proces kompilacji w oparciu o platformę ROCm. Natomiast ścieżka kompilacji **dla urządzeń NVIDIA ** korzysta ze standardowego kompilatora CUDA (CUDA > nvcc). W tym przypadku, jako implementacja HIP zostaje użyty wrapper tłumaczący poszczególne wywołania API HIP na wywołania API CUDA. Narzut wydajnościowy takiego rozwiązania jest znikomy. Dzięki temu możliwe staje się skompilowanie kodu HIP przy użyciu kompilatora NVIDIA, tak jakby był to kod CUDA.

HIP umożliwia również korzystanie z funkcjonalności specyficznych dla architektur NVIDIA oraz AMD. Muszą jednak one być ujęte w dyrektywy warunkowe preprocesora. Co więcej, możliwe jest mieszanie kodu CUDA z kodem HIP.

Platforma ROCm

Patrząc na rozwiązania firmy NVIDIA, CUDA to nie tylko model programistyczny ale również platforma zawierające zbiór narzędzi, sterowników, bibliotek, a więc całe środowisko do pracy z kartami GPU firmy NVIDIA. Odpowiednikiem platformy CUDA po stronie AMD jest ROCm. Zawiera ona zbiór rozwiązań do programowania kart GPU firmy AMD.

W powyższym kontekście, HIP jest składową platformy ROCm, i dostarcza wysokopoziomowy interfejs do programowania GPU. Integruje się on z niżej poziomowymi rozwiązaniami firmy AMD, a jednocześnie umożliwia integrację ze środowiskiem CUDA.

Przykładem rozróżnienia na HIP oraz ROCm są biblioteki obliczeniowe dostarczane przez AMD w ramach pakietu ROCm. Niektóre z nich występują w dwóch wariantach: z przedrostkiem "hip" oraz z przedrostkiem " roc", na przykład hipBLAS/rocBLAS.

  • Biblioteki rocXXX to zoptymalizowane implementacje przeznaczone tylko dla urządzen AMD. Są one odpowiednikami analogicznych bibliotek firmy NVIDIA np. rocBLAS vs cuBLAS.
  • Biblioteki hipXXX dostarczają warstwę abstrakcji pozwalającą na przenośność kodu między urządzeniami AMD oraz NVIDIA. Wymagają one do działania (kompilacji) obecności biblioteki zawierającej faktyczną implementację na daną platformę. Na przykładzie hipBLAS będą to: rocBLAS – dla urządzeń AMD, oraz cuBLAS – dla urządzeń NVIDIA.

Terminologia

Terminologia używana w HIP jest zgodna z CUDA – patrz porównanie składni CUDA, HIP, OpenCL. Szerszy opis jest dostępny w sekcji CUDA > terminologia. Należy pamiętać, że ⚠️ w architekturze AMD warp (wavefront) liczy 64 wątki, oraz że niektóre funkcjonalności CUDA nie są dostępne w HIP.

Bezpośredni dostęp do GPU

Akceleratory AMD umożliwiają korzystanie z rozwiązań bezpośredniego dostępu do danych takich jak DMA (direct memory access) czy RDMA (remote direct memory access). Pozwalają ona na dostęp do danych z pamięci GPU przez inne urządzenie w ramach tego samego komputera, lub sieci, bez interakcji z CPU. Odpowiada to technologii CUDA > GPUDirect.

AMD Infinity Fabric

AMD opracowało również technologię Infinity Fabric, która jest odpowiednikiem NVLink i służy do przyśpieszenia bezpośredniej komunikacji między akceleratorami GPU w ramach jednego komputera. Dodatkowo, w systemach bazujących na procesorach AMD, Infinity Fabric wykorzystywany jest również do komunikacji CPU-CPU oraz CPU-GPU.

W HIP wsparcie dla komunikacji bezpośredniej jest w fazie eksperymentalnej (patrz moduł PeerToPeer Device Memory Access).

Informacje o standardzie

HIP to stosunkowo młoda technologia, rozwijana przez firmę AMD od ok. 2016 roku. Jest podstawowym modelem programistycznym dla urządzeń GPU AMD, wchodzącym w skład szerokiego ekosystemu AMD ROCm. HIP jest mniej dojrzały od konkurencyjnej CUDA, która ma sobą znacznie dłuższą ścieżkę rozwoju. Przede wszystkim dostarczane narzędzia, a także dokumentacja, są mniej dopracowane i mniej zaawansowane. HIP oraz pakiet ROCm są jednak nieustannie rozwijane, a braki sukcesywnie zmniejszane.

Przez wiele lat na rynku kart GPU wykorzystywanych do obliczeń ogólnego przeznaczenia GPGPU (general-purpose computing on graphics processing units) NVIDIA była bezkonkurencyjna. W efekcie model CUDA stał się bardzo powszechny i wszystkie rozwiązania GPGPU były skoncentrowane wokół niego. Obecnie, dzięki rosnącej konkurencji pomiędzy NVIDIA a AMD w tym segmencie GPU, platforma ROCm oraz HIP zyskują na znaczeniu. Wyrazem tego jest m.in. wybór akceleratorów GPU AMD do budowy najwydajniejszych superkomputerów na świecie (na przykładzie listy TOP500 z czerwca 2023 są to: Frontier – 1 miejsce, oraz LUMI – 3 miejsce). W ślad za rozwojem sprzętowym AMD postępuje dostosowywanie głównych pakietów obliczeniowych do środowiska ROCm.

Wsparcie frameworków AI dla GPU AMD

Najpopularniejsze biblioteki i frameworki uczenia maszynowego (TensorFlow, PyTorch) pierwotnie powstały tylko ze wsparciem dla GPU NVIDIA (CUDA). W ostatnim czasie uległo to zmianie i obecnie są również dostępne backendy wykorzystujące HIP. Dzięki temu możliwe jest korzystanie z frameworków AI także na akceleratorach AMD.

Należy jednak dodać, że HIP/ROCm koncentrują się na wyspecjalizowanych kartach GPU, przeznaczonych do centrów obliczeniowych oraz stacji roboczych (serie AMD Instinct oraz Radeon Pro). Większość układów GPU przeznaczonych do urządzeń konsumenckich nie jest oficjalnie wspierana w najnowszych pakietach (patrz lista wspieranych GPU). W tym zakresie rozwiązania AMD mają ograniczoną dostępność i odstają od NVIDIA, gdzie właściwie każda standardowa karta pozwala na programowanie z wykorzystaniem CUDA.

Pakiet ROCm

Do tworzenia oraz uruchamiania aplikacji korzystających z HIP niezbędny jest pakiet oprogramowania i narzędzi ROCm. W jego skład wchodzą między innymi:

  • kompilatory hipcc, amdclang++,
  • biblioteki uruchomieniowe HIP,
  • podstawowe narzędzia (profiler, debugger),
  • biblioteki programistyczne (np. rocBLAS/hipBLAS, rocRAND/hipRAND).

W porównaniu do rozwiązań NVIDIA, pakiet ROCm to pojedynczy odpowiednik pakietów CUDA Toolkit oraz HPC SDK. Jednakże posiada on mniej rozwiązań dotyczących obliczeń rozproszonych – nie zawiera dystrybucji MPI oraz ScaLAPACK, a biblioteki obliczeniowe nie posiadają wariantów do liczenia na wielu kartach GPU. Z tego typu rozwiązań dostępna jest tylko biblioteka RCCL do komunikacji między wieloma GPU.

Aktualna wersja pakietu to AMD ROCm 5.6.0, wydana w maju 2023. Zawartość można obejrzeć sprawdzając dokumentację składowych pakietu. W przypadku bibliotek obliczeniowych, informacje o numerach wersji można znaleźć w informacjach o wydaniu. W przeciwieństwie do rozwiązań NVIDIA, składowe pakietu ROCm są rozwijane jako oprogramowanie open-source. Kod źródłowy udostępniany jest w serwisie GitHub w ramach projektów ROCm Core Technology oraz ROCm Software Platform.

Pakiet ROCm jest na bieżąco rozwijany i często aktualizowany. Wersje major (np. 5.0) wiążą się zwykle z większymi zmianami i wydawane są w odstępach mniej więcej od 1 do 1.5 roku. Uaktualnienia (wersje minor, np. 5.4, 5.5) ukazują się kilka razy w roku. Dodatkowo mogą pojawiać się także wersje patch (np. 5.5.1) zawierające poprawki błędów danej wersji.

Kompilatory

Tak jak w CUDA, aby skompilować program HIP potrzebny jest specjalny kompilator. Jego rolą jest identyfikacja kodu hosta oraz kodu akcelerowanego i kompilacja jednego pod CPU, a drugiego pod wybraną architekturę GPU. Ogólnym kompilatorem przewidzianym do pracy zarówno z GPU AMD jak i GPU NVIDIA jest hipcc. Drugi kompilator z pakietu ROCm (amdclang++) dotyczy stricte kompilacji pod architekturę AMD.

hipcc

Podstawowy, wchodzący w skład pakietu ROCm kompilator HIP. Umożliwia zarówno kompilację pod akceleratory AMD jak i NVIDIA. Ściśle mówiąc hipcc to wrapper, który w zależności od wybranej platformy GPU przekazuje zadanie kompilacji do odpowiedniego kompilatora:

  • w przypadku GPU NVIDIA będzie to CUDA > nvcc,
  • w przypadku GPU AMD będzie to amdclang++ (a dokładniej kompilator ROCmCC).

W szczególności opcje kompilacji są zależne od docelowego kompilatora. Rola hipcc ogranicza się tylko do ustawienia podstawowych zmiennych dla danego środowiska (np. ścieżki do plików nagłówkowych i linkowania). Dla wygody użytkownika, hipcc traktuje wszystkie pliki źródłowe (niezależnie od rozszerzenia) jako zawierające kod HIP.

Więcej informacji

amdclang++ (ROCmCC)

ROCmCC to właściwy kompilator przewidziany dla architektury GPU AMD, oparty o platformę LLVM i kompilator clang (co tłumaczy potoczną nazwę HIP-Clang). Jest on rozwijany przez AMD jako otwartoźródłowy fork projektu LLVM. Jest samodzielnym kompilatorem dla języka C++, który dodatkowo potrafi kompilować kod HIP. Oprócz tego umożliwia wykorzystanie kart GPU AMD przy użyciu technologii OpenCL oraz OpenMP (offloading). W porównaniu do rozwiązań CUDA jest to mniej więcej odpowiednik nvcc połączonego z nvc++ (patrz CUDA > kompilatory.

Kompilator ROCmCC jest elementem pakietu ROCm. Wywołuje się go poprzez komendę amdclang++, która posiada interfejs standardowego kompilatora clang++ (te same opcje kompilacji). Pozwala na szczegółową kontrolę procesu kompilacji. Do kompilacji kodu HIP wymaga plików z rozszerzeniem .hip lub flagi -x hip.

W kontekście kompilacji pod GPU AMD, hipcc może być postrzegany jako drugi interfejs tego samego kompilatora. W dokumentacji są opisane różnice między hipcc a amdclang.

Optymalizacje dla CPU AMD

Standardowy ROCmCC (pakiet rocm-llvm) może zostać rozszerzony poprzez dodatkowy pakiet rocm-llvm-alt o optymalizacje specyficzne dla CPU firmy AMD. Optymalizacje te wywodzą się z zoptymalizowanego przez AMD pod swoją architekturę kompilatora AOCC.

Główna opcja do wskazywania architektury GPU dla której ma zostać skompilowany kod akcelerowany to --offload-arch=<gpu>. Parametr ten można podawać kilkukrotnie, specyfikując wiele architektur. Plik binarny wygenerowany przez kompilator będzie kompatybilny tylko z wyspecyfikowanymi architekturami.

Architektury GPU AMD i szczegóły postaci binarnej

⚠️ AMD nie posiada odpowiednika PTX oraz pojęcia compute-capability z CUDA.

W okresie wprowadzania do użytku środowiska ROCm karty GPU firmy AMD były oparte o architekturę GCN (Graphics Core Next). W związku z tym najczęściej w odniesieniu do kodu binarnego (lub jego asemblerowego zapisu) mówi się o AMD GCN assembly (mimo tego, że najnowsze architektury AMD noszą inne nazwy).

Obecne architektury stosowane w GPU AMD to RDNA oraz CDNA. Pierwsza zoptymalizowana jest pod kątem generowania grafiki na bieżąco, druga pod kątem akceleracji obliczeń. Stąd właśnie seria kart AMD Instinct przewidziana dla segmentu HPC jest oparta o architekturę CDNA.

Każda z architektur ma swoje warianty. Kompilując kod po konkretną kartę GPU należy sprawdzić nazwę kodową wariantu architektury danego GPU:

Jak można zauważyć, między różnymi GPU zachodzi kompatybilność na poziomie binarnym. Przykładowo gfx90a to architektura zarówno dla MI210 jak i MI250; kod skompilowany w tej architekturze będzie można uruchomić na obydwóch kartach.

Więcej informacji

clang

Otwartoźródłowy kompilator języków C/C++ rozwijany w ramach projektu LLVM. Potrafi również kompilować kod CUDA (stanowi alternatywę dla nvcc). Wybrany przez AMD na podstawę swojego kompilatora dla środowiska ROCm. Zmiany wprowadzane przez AMD do platformy LLVM pod kątem obsługi własnych GPU, są stopniowo migrowane do głównych gałęzi projektu. Dzięki temu możliwe jest używanie standardowego kompilatora clang do kompilacji kodu HIP na GPU AMD. Warto mieć na uwadze, że nie każda instalacja clang może mieć włączone wsparcie dla backendu AMD (podobnie jak nie każda ma wsparcie dla CUDA). W takim przypadku konieczne jest samodzielne zbudowanie kompilatora z odpowiednimi flagami.

Dużo opcji clang, które do tej pory dotyczyły tylko CUDA otrzymało alternatywne nazwy z "gpu" w miejsce "cuda" w nazwie parametru. Przykładowo -fgpu-rdc odpowiada -fcuda-rdc. Główna komenda do wskazywania architektur to --offload-arch (w miejsce --cuda-gpu-arch).

Obecna wersja stabilna to 16.0.6, wydana w czerwcu 2023.

Więcej informacji

Narzędzia

AMD udostępnia podstawowy zestaw narzędzi (profiler, debugger) do pracy z aplikacjami HIP/ROCm. Są to rozwiązania mniej wygodne w użytkowaniu niż ich odpowiedniki znane z CUDA, jednak dostarczają one pełny zestaw istotnych funkcjonalności. Oprócz tego dostępne jest narzędzie ułatwiające migrację z CUDA do HIP.

  • HIPify ─ narzędzie wchodzące w skład pakietu AMD ROCm umożliwiające automatyczne lub półautomatyczne tłumaczenie kodu CUDA do HIP.
  • rocprof ─ profiler aplikacji korzystających z HIP/ROCm. Udostępnia tylko interfejs wiersza poleceń (CLI), a wejście i wyście obsługiwane są poprzez pliki tekstowe. Dostarcza informacji na temat przebiegu działania aplikacji oraz wykorzystania zasobów sprzętowych w czasie.
  • ROCgdb ─ debugger aplikacji wykorzystujących HIP oraz platformę ROCm. Bazuje na GNU gdb. Umożliwia wykrywanie błędów, przerywanie działania programu, wykonywanie kodu krok po kroku, podglądanie wartości zmiennych, stosu wywołań funkcji, itp.

Linki

Ostatnia aktualizacja: 15 września 2023