Przejdź do treści

OpenCL

logo OpenCL

Wprowadzenie

OpenCL (Open Computing Language) to otwarty standard określający model programowania dla platform akcelerowanych, który powstał jako alternatywa dla CUDA. Jest rozwiązaniem nieograniczającym się tylko do kart GPU, tylko jednego producenta. Udostępnia jednolity interfejs pozwalający programować różne typy akceleratorów, zaliczając do tego grona również wielordzeniowe CPU.

Początkowo stanowił najpoważniejszą alternatywę dla CUDA do programowania akceleratorów GPU. Przez pewien czas był podstawową metodą programowania akceleratorów GPU AMD. Wraz z rozwojem nowszych modeli programowania wspierających architektury GPU, takich jak HIP czy wysokopoziomowy SYCL, rola OpenCL na tym polu zmalała. Obecnie stosowany jest coraz częściej do programowania mniej rozpowszechnionych, wąsko wyspecjalizowanych układów obliczeniowych.

Dostępność

Podstawowym językiem programowania do pracy z OpenCL jest język C. Interakcja z akceleratorem jest realizowana poprzez funkcje biblioteczne zdefiniowane przez standard OpenCL (API w języku C) oraz funkcje akcelerowane (wyrażane w nieznacznie zmodyfikowanym języku C). Dostępne są również oficjalne rozwiązania ułatwiające pracę z językiem C++.

Program OpenCL kompiluje się przy użyciu standardowego kompilatora C/C++, jednakże niezbędne jest posiadanie zainstalowanej tzw. implementacji OpenCL. Rolą implementacji jest dostarczenie biblioteki OpenCL oraz sterownika potrafiącego kompilować kod akcelerowany pod wspierane przez daną implementację rodzaje akceleratorów (więcej w sekcji implementacje).

Szczegóły

OpenCL realizuje klasyczny model z podziałem na kod hosta (wykonywany na CPU) i kod akceleratora w postaci odrębnych funkcji (tzw. kerneli). W odróżnieniu od innych modeli programowania akceleratorów (CUDA, HIP, SYCL) wymaga aby kod kerneli był wyodrębniony do osobnych plików. Ewentualnie kernele mogą być deklarowane w kodzie hosta jako literał tekstowy (ciąg znaków char *).

Model programistyczny

model programistyczny OpenCL z separacją kodu akcelerowanego [źródło]

Kod hosta poprzez funkcje biblioteczne kontroluje zachowanie akceleratora i zleca mu do wykonania kernele. Kernele wyraża się w tzw. języku OpenCL C, czyli języku C z drobnymi modyfikacjami składniowymi. Oficjalny standard wspiera również pracę w języku C++, poprzez wrappery C++ ułatwiające korzystanie z biblioteki OpenCL oraz język C++ for OpenCL do pisania kerneli. C++ for OpenCL rozszerza OpenCL C o możliwość korzystania z większości funkcjonalności języka C++17 w kodzie akcelerowanym.

C++ for OpenCL

język "C++ for OpenCL" jako rozszerzenie OpenCL C na C++ [źródło]

OpenCL a inne języki programowania

Dzięki separacji kodu akcelerowanego od kodu hosta (i ich osobnej kompilacji), możliwe jest stosunkowo łatwe użycie OpenCL z poziomu innych języków programowania. W tym celu wystarczy wołać funkcje z biblioteki OpenCL (API C) z poziomu innego języka. W przypadku języków dobrze współpracujących z C (jak np. Fortran) można to zrobić bezpośrednio, Istnieją również (nieoficjalne) biblioteki pośredniczące dla innych języków.

Należy jednak mieć na uwadze, że w każdym przypadku część akcelerowana musi zostać dostarczona jako kod OpenCL C.

W porównaniu do innych modeli programowania akceleratorów, OpenCL może być postrzegany jako **niskopoziomowe rozwiązanie ** bazujące na języku C. Przejawia się to w tym, że wiele podstawowych czynności, które w wyżej poziomowych modelach programowania mogą zostać wykonane automatycznie, w OpenCL trzeba dokonywać manualnie. Zwiększa to poziom kontroli nad zachowaniem sprzętu, jednakże powoduje, że nawet prosty program staje się stosunkowo złożony.

Schemat programu

Przykład kodu OpenCL

przykład programu OpenCL z kernelem przekazywanym w formie tekstowej [źródło]

Typowy program OpenCL (w języku C) ma mniej więcej następującą strukturę, a każdej akcji odpowiada wywołanie jednej lub kilku funkcji.

Czynności początkowe:

  • wybór akceleratora,
  • utworzenie powiązanego z nim tzw. obiektu kontekstu (cl_context) – clCreateContext
  • utworzenie kolejki do zlecania obliczeń – clCreateCommandQueue

Kolejnym krokiem jest przygotowanie kernela do wykonania na akceleratorze, co wykonuje się według schematu:

  • wczytanie tzw. programu, czyli kodu źródłowego z kernelami,
  • utworzenie obiektu programu i zlecenie kompilacji kerneli – funkcje z grupy clCreateProgramWith... oraz clBuildProgram
  • wybranie odpowiedniego kernela z programu i finalnie utworzenie obiektu kernela gotowego do wykonywania obliczeń – clCreateKernel

Dopiero tak przygotowany kernel można uruchomić na akceleratorze. Wywołanie kernela, wymaga podania argumentów kernela (clSetKernelArg) oraz konfiguracji poziomu równoległości, co dokonuje się w równie niskopoziomowy sposób. Oczywiście wcześniej trzeba przygotować pamięć dla akceleratora. Alokacja pamięci w przestrzeni adresowej akceleratora odbywa się poprzez utworzenie powiązanych z odpowiednim kontekstem buforów (clCreateBuffer). Transfery danych oraz wywołania kernela zleca się korzystając z obiektu kolejki (funkcje clEnqueueWriteBuffer, clEnqueueReadBuffer oraz clEnqueueNDRangeKernel).

Po zakończeniu pracy wszystkie używane wcześniej obiekty należy ręcznie zdealokować odpowiednimi funkcjami (clRelease...).

Kompilacja kerneli

Patrząc na opisany schemat użycia, widać, że kompilacja kodu akcelerowanego odbywa się dopiero w czasie wykonania programu (runtime). Faktyczna kompilacja wykonywana jest przez tzw. driver (sterownik) OpenCL obsługujący docelowe urządzenie/akcelerator. Dzięki temu, że obiekt programu jest skojarzony z kontekstem, który to z kolei jest przypisany do konkretnego urządzenia, biblioteka OpenCL jest w stanie przekazać zadanie kompilacji do odpowiedniego sterownika. Takie podejście pozwala na kompilację tego samego kodu akcelerowanego pod różne akceleratory i stwarza możliwość wykorzystywania przez jeden program akceleratorów różnego typu.

Na wielu platformach możliwa jest także prekompilacja kerneli do odpowiedniej postaci pośredniej lub binarnej – potrzeba jednak do tego odpowiedniego kompilatora. W takiej sytuacji, w programie użytkownika wystarczy wczytać plik ze skompilowaną wersją kernela. Jest to rozwiązanie szczególnie często stosowane w sytuacjach gdy kompilacja kerneli zajmuje dużo czasu (na przykład przy pracy z układami FPGA). Ogranicza jednak ono przenośność kerneli tylko do platform o kompatybilnym formacie binarnym lub pośrednim.

OpenCL a CUDA

Jako rozwiązanie niżej poziomowe, OpenCL jest w użyciu pod wieloma względami bardziej zbliżony do wykorzystania CUDA za pośrednictwem driver API niż standardowego runtime API (patrz CUDA > Biblioteka CUDA). Przykładowo, w CUDA driver API wymaga kernele także są wczytywane z zewnętrznych plików – w tym wypadku jednak zawsze w postaci prekompilowanej.

Pomimo innego podejścia i własnej terminologii, większość konceptów stosowanych w OpenCL znajduje swoje odpowiedniki w CUDA (patrz spis najistotniejszych analogii oraz CUDA > Terminologia).

OpenCL a SYCL

Modele programistyczne OpenCL oraz SYCL na poziomie podstawowych koncepcji i terminologii są do siebie zbliżone. Podobnie stosuje się kolejki do delegowania obliczeń, zarządzanie pamięcią odbywa się poprzez obiekty buforów, a siatkę obliczeń definiuję się za pomocą zakresów (NDRange). W OpenCL nie występują jednak pewne koncepty rozwinięte w SYCL. Przykładem mogą być akcesory, przez co zarządzanie transferami pamięci w OpenCL musi odbywać się ręcznie.

Podobieństwo modeli wynika z tego, że SYCL powstał w ramach prac nad standardem OpenCL jako jego wysokopoziomowy interfejs. Początkowo był on nakładką na OpenCL, której rolą było ułatwienie programowania kerneli – m.in. poprzez możliwość mieszania kodu hosta z kodem akcelerowanym oraz dostarczenie nowoczesnych funkcjonalności języka C++. Obecnie SYCL jest rozwijany jako osobny standard i w znacznej mierze przejmuje rolę uniwersalnego sposobu programowania akceleratorów różnego typu (przede wszystkim GPU oraz wielordzeniowe CPU).

Informacje o standardzie

Pierwsza wersja specyfikacji została udostępniona w 2009 roku. Standard jest systematycznie rozwijany przez Khronos Group (konsorcjum koordynujące rozwój wielu otwartych standardów). Obecna główna wersja to OpenCL 3.0, opublikowana we wrześniu 2020 roku, z późniejszymi drobnymi poprawkami. Aktualna wersja to 3.0.14, wydana w kwietniu 2023 roku.

Wersja 3.0 uporządkowała standard, dokonując unifikacji wersji 1.2 oraz 2.x. Dzieli ona funkcjonalności OpenCL na wymagane oraz opcjonalne. Część wymagana do zgodności ze standardem pokrywa się z OpenCL 1.2, z kolei funkcjonalności wprowadzone do standardu w wersjach 2.x uznano w wersji 3.0 za opcjonalne. Takie rozwiązanie z jednej strony ogranicza konkurencyjność OpenCL względem innych modeli programowania akceleratorów, gdyż zgodność ze standardem nie gwarantuje obsługi nowych funkcjonalności. Jednocześnie ułatwia to implementację standardu dla wąsko wyspecjalizowanych układów obliczeniowych, na których zaimplementowanie wysokopoziomowych rozwiązań jest często niemożliwe lub nieefektywne. Przykładem takiej funkcjonalności może być pamięć zunifikowana (unified memory). Duża opcjonalność na poziomie standardu jest krokiem w kierunku ugruntowania OpenCL jako niskopoziomowego modelu do programowania wąsko wyspecjalizowanych platform akcelerowanych.

Język programowania kerneli "C++ for OpenCL" jest oparty o standard OpenCL 3.0 oraz C++17. Posiada on swoją własną numerację, która odpowiada unifikacji standardu OpenCL.

  • wersja 1.0, kompatybilna wstecznie z OpenCL 2.0
  • wersja 2021, kompatybilna wstecznie z OpenCL 3.0

Implementacje

Kod OpenCL może być uruchamiany na wielu platformach sprzętowych. Pewne rozeznanie, które sprzęty posiadają wsparcie dla OpenCL, dostarcza lista produktów zgodnych ze standardem OpenCL, utrzymywana przez grupę Khronos. Lista ta obejmuje tylko te implementacje i urządzenia, które przeszły w całości oficjalne testy zgodności z pełną wersją danego standardu; szukając implementacji pod kątem konkretnego sprzętu nie należy ograniczać się tylko do tej listy.

Implementacje OpenCL

historia implementacji zgodnych ze standardem – dane do roku 2017 [źródło]

Najważniejsi producenci układów CPU oraz GPU (AMD, Intel oraz NVIDIA) dostarczają implementacje OpenCL dla swoich architektur. Najczęściej wchodzą one w skład podstawowych pakietów narzędzi do tworzenia oprogramowania na te platformy (takich jak Intel oneAPI Base Toolkit, CUDA Toolkit, ROCm). W niektórych przypadkach możliwe jest również osobne pobranie środowiska do uruchamiania OpenCL (np. runtime dla Intel). Istnieją także alternatywne, otwartoźródłowe implementacje, takie jak PoCL (Portable Computing Language)

Wsparcie OpenCL jest powszechne w segmencie mobilnych procesorów o architekturze Arm oraz towarzyszących im układów GPU (Samsung, Qualcomm, Google). Dodatkowo, OpenCL jest często stosowany jako jedna z metod programowania wyspecjalizowanych oraz rozwijających się architektur takich jak układy FPGA (AMD/Xilinx, Intel), DSP (Texas Instruments) czy RISC-V.

Linki

Ostatnia aktualizacja: 15 września 2023