Nowa technologia jest jak nowo powstały gatunek ewolucyjny. Dziwne stworzenie, w przeciwieństwie do wielu staruszków. Czasami niezręcznie, czasami zabawnie. I z początku jego nowe cechy nie wydają się być odpowiednie dla tego nadającego się do zamieszkania i stabilnego świata.

Mija jednak trochę czasu i okazuje się, że początkujący biega szybciej, skacze wyżej i ogólnie mocniejszy. I zjada więcej much niż jego sąsiedzi w retrogradacji. I wtedy ci sami sąsiedzi zaczynają rozumieć, że nie warto się kłócić z tą dawną niezdarą. Lepiej się z nim zaprzyjaźnić, a jeszcze lepiej zorganizować symbiozę. Patrzysz, a much będzie więcej.

Technologia GPGPU (General-Purpose Graphics Processing Units) od dawna istnieje tylko w teoretycznych obliczeniach bystrych naukowców. Jak inaczej? Zaproponuj radykalną zmianę w minionych dziesięcioleciach proces obliczeniowy, powierzając obliczenia jego równoległych gałęzi karcie graficznej - tylko teoretycy są do tego zdolni.

Logo technologii CUDA przypomina nam, że wyrosła w głębinach
Grafika 3D.

Ale technologia GPGPU nie zamierzała długo gromadzić kurzu na łamach uniwersyteckich czasopism. Wytrzebiwszy pióra swoich najlepszych cech, przyciągnęła uwagę producentów. Tak narodziła się CUDA, implementacja GPGPU na procesorach graficznych NVIDIA GeForce.

Dzięki CUDA technologie GPGPU stały się głównym nurtem. A teraz tylko najbardziej krótkowzroczny i najbardziej leniwy deweloper systemów programowania nie deklaruje wsparcia dla CUDA swoim produktem. Publikacje informatyczne miały zaszczyt przedstawić szczegóły tej technologii w licznych obszernych artykułach popularnonaukowych, a konkurenci pilnie zasiedli do szablonów i kompilatorów krzyżowych, aby opracować coś podobnego.

Uznanie społeczne to nie tylko marzenie aspirujących gwiazd, ale także nowo powstających technologii. A CUDA miała szczęście. Jest słyszana, rozmawiają o niej i piszą o niej.

Po prostu piszą tak, jakby nadal dyskutowali o GPGPU w grubych czasopismach naukowych. Bombardują czytelnika mnóstwem terminów, takich jak „grid”, „SIMD”, „warp”, „host”, „pamięć strukturalna i stała”. Zanurzają go na sam szczyt schematów organizacji GPU nVidii, prowadzą przez kręte ścieżki równoległych algorytmów i (najsilniejszy ruch) pokazują długie listy kodu w języku C. W efekcie okazuje się, że na wejściu do artykułu mamy świeżą i palącą chęć zrozumienia czytnika CUDA, a na wyjściu - ten sam czytnik, ale z opuchniętą głową wypełnioną bałaganem faktów, schematów, kodu , algorytmy i terminy.

Tymczasem celem każdej technologii jest ułatwienie nam życia. A CUDA świetnie sobie z tym radzi. Wyniki jej pracy przekonają każdego sceptyka lepiej niż setki schematów i algorytmów.

Daleko nigdzie

CUDA wspierana przez superkomputery o wysokiej wydajności
NVIDIA Tesla.

A jednak, zanim przyjrzymy się wynikom prac CUDA w zakresie ułatwiania życia przeciętnemu użytkownikowi, warto zrozumieć wszystkie jego ograniczenia. Tak jak z dżinem: każde pragnienie, ale jedno. CUDA ma również swoje pięty achillesowe. Jednym z nich są ograniczenia platform, na których może pracować.

Lista kart graficznych wyprodukowanych przez nVidia, które obsługują CUDA, jest przedstawiona na specjalnej liście o nazwie Produkty z obsługą CUDA. Lista jest imponująca, ale łatwo ją sklasyfikować. Nie odmawia się wsparcia CUDA:

Modele nVidia GeForce z serii 8, 9, 100, 200 i 400 z minimum 256 MB pamięci wideo na pokładzie. Wsparcie obejmuje zarówno karty stacjonarne, jak i mobilne.

Zdecydowana większość stacjonarnych i mobilnych kart graficznych to nVidia Quadro.

Wszystkie rozwiązania serii netbooków nvidia ION.

Wysokowydajne rozwiązania superkomputerowe HPC (High Performance Computing) i nVidia Tesla wykorzystywane zarówno do komputerów osobistych, jak i do organizacji skalowalnych systemów klastrowych.

Dlatego przed użyciem oprogramowania opartego na CUDA warto sprawdzić tę listę ulubionych.

Oprócz samej karty graficznej obsługa CUDA wymaga odpowiedniego sterownika. To on jest łącznikiem między procesorem centralnym a procesorem graficznym, działając jako swego rodzaju interfejs programowy umożliwiający dostęp do kodu i danych programu do wielordzeniowego skarbca GPU. Aby mieć pewność, że nie popełnisz błędu, nVidia zaleca odwiedzenie strony ze sterownikami i pobranie najnowszej wersji.

...ale proces

Jak działa CUDA? Jak wytłumaczyć złożony proces obliczeń równoległych na określonej architekturze sprzętowej GPU bez zanurzania czytelnika w otchłań konkretnych terminów?

Możesz spróbować to zrobić, wyobrażając sobie, jak centralny procesor wykonuje program w symbiozie z procesorem graficznym.

Architektonicznie jednostka centralna (CPU) i jej graficzny odpowiednik (GPU) są rozmieszczone inaczej. Jeśli narysujemy analogię ze światem motoryzacji, to procesorem jest kombi, jeden z tych, które nazywa się „stodoła”. Wygląda jak samochód osobowy, ale jednocześnie (z punktu widzenia twórców) „i Szwajcar, i żniwiarz, i gracz na fajce”. Pełni jednocześnie rolę małej ciężarówki, autobusu i przerośniętego hatchbacka. Krótko mówiąc, uniwersalny. Ma niewiele rdzeni cylindrów, ale „wyciągają” prawie każde zadanie, a imponująca pamięć podręczna może pomieścić mnóstwo danych.

Ale GPU to samochód sportowy. Jest tylko jedna funkcja: jak najszybciej dostarczyć pilota do mety. Dlatego nie ma dużego bagażnika pamięci, żadnych dodatkowych miejsc. Ale jest setki razy więcej cylindrów rdzeniowych niż procesor.

Dzięki CUDA programiści GPGPU nie muszą rozumieć złożoności programowania.
dla silników graficznych, takich jak DirectX i OpenGL

W przeciwieństwie do procesora centralnego, który jest w stanie rozwiązać każde zadanie, w tym grafikę, ale ze średnią wydajnością, procesor graficzny jest przystosowany do szybkiego rozwiązania jednego zadania: zamiany stosów wielokątów na wejściu w kilka pikseli na wyjście. Co więcej, zadanie to można rozwiązać równolegle na setkach stosunkowo prostych rdzeni obliczeniowych w ramach GPU.

Czym więc może być tandem kombi i samochodu sportowego? Praca CUDA przebiega mniej więcej tak: program jest wykonywany na procesorze, dopóki nie zawiera fragmentu kodu, który można wykonać równolegle. Następnie, zamiast powoli uruchamiać się na dwóch (nawet ośmiu) rdzeniach najfajniejszego procesora, jest przenoszony na setki rdzeni GPU. Jednocześnie skraca się kilkakrotnie czas wykonania tej sekcji, co oznacza, że ​​skraca się również czas wykonania całego programu.

Technologicznie dla programisty nic się nie zmienia. Kod programów CUDA napisany jest w języku C. Dokładniej, w swoim specjalnym dialekcie „C ze strumieniami” (C ze strumieniami). Opracowane w Stanford rozszerzenie języka C nazywa się Brook. Interfejs, który przesyła kod Brooka do GPU, jest sterownikiem karty graficznej obsługującej CUDA. Organizuje cały proces przetwarzania tej sekcji programu w taki sposób, że GPU wygląda dla programisty jak koprocesor CPU. Podobnie jak używanie koprocesora matematycznego we wczesnych dniach komputerów osobistych. Wraz z pojawieniem się Brooka, kart graficznych z obsługą CUDA i sterowników do nich, każdy programista uzyskał dostęp do GPU w swoich programach. Ale wcześniej ten szamanizm należał do wąskiego kręgu elity, która przez lata doskonaliła techniki programowania dla silników graficznych DirectX lub OpenGL.

W beczce tego pretensjonalnego miodu – chwali się CUDA – warto włożyć muchę w maść, czyli ograniczenia. Nie każde zadanie, które trzeba zaprogramować, nadaje się do rozwiązania za pomocą CUDA. Osiągnięcie przyspieszenia w rozwiązywaniu rutynowych zadań biurowych nie będzie możliwe, ale możesz zaufać CUDA, aby obliczyć zachowanie tysięcy myśliwców tego samego typu w World of Warcraft. Ale to zadanie wyssane z palca. Rozważmy przykłady tego, co już CUDA rozwiązuje bardzo skutecznie.

Sprawiedliwe dzieła

CUDA to bardzo pragmatyczna technologia. Po zaimplementowaniu jego obsługi w swoich kartach graficznych, nVidia całkiem słusznie spodziewała się, że baner CUDA zostanie podchwycony przez wielu entuzjastów zarówno w środowisku uniwersyteckim, jak iw handlu. I tak się stało. Projekty oparte na CUDA żyją i mają się dobrze.

NVIDIA PhysX

Reklamując kolejne arcydzieło gier, producenci często podkreślają jego realizm 3D. Ale bez względu na to, jak prawdziwy jest świat gry 3D, jeśli podstawowe prawa fizyki, takie jak grawitacja, tarcie, hydrodynamika, zostaną niepoprawnie zaimplementowane, fałsz będzie natychmiast odczuwalny.

Jedną z funkcji silnika fizyki NVIDIA PhysX jest realistyczna praca z tkankami.

Implementacja algorytmów komputerowej symulacji podstawowych praw fizycznych jest bardzo pracochłonnym zadaniem. Najbardziej znane firmy w tej dziedzinie to irlandzka firma Havok ze swoją wieloplatformową fizyką Havok Physics oraz Californian Ageia - protoplasta pierwszego na świecie procesora fizycznego (PPU - Physics Processing Unit) i odpowiadającego mu silnika fizyki PhysX. Pierwszy z nich, choć przejęty przez Intela, obecnie aktywnie działa w dziedzinie optymalizacji silnika Havok dla kart graficznych ATI i procesorów AMD. Ale Ageia ze swoim silnikiem PhysX stała się częścią nVidii. W tym samym czasie nVidia rozwiązała dość trudne zadanie dostosowania PhysX do technologii CUDA.

Jest to możliwe dzięki statystykom. Udowodniono statystycznie, że bez względu na to, jak złożone jest renderowanie GPU, niektóre z jego rdzeni są nadal bezczynne. To na tych rdzeniach działa silnik PhysX.

Dzięki CUDA lwia część obliczeń związanych z fizyką świata gry zaczęła być wykonywana na karcie graficznej. Uwolniona moc centralnego procesora została rzucona na inne zadania rozgrywki. Na wynik nie trzeba było długo czekać. Według ekspertów wzrost wydajności rozgrywki z PhysX działającym na CUDA wzrósł przynajmniej o rząd wielkości. Wzrosła również prawdopodobieństwo realizacji praw fizycznych. CUDA zajmuje się rutynowymi obliczeniami implementacji tarcia, grawitacji i innych znanych nam rzeczy dla obiektów wielowymiarowych. Teraz nie tylko bohaterowie i ich ekwipunek doskonale wpisują się w znane nam prawa świata fizycznego, ale także kurz, mgła, fala uderzeniowa, płomień i woda.

Wersja CUDA pakietu kompresji tekstur NVIDIA Texture Tools 2

Lubisz realistyczne obiekty w nowoczesnych grach? Warto to podziękować twórcom tekstur. Ale im więcej rzeczywistości w fakturze, tym większa jest jej objętość. Im więcej zabiera cennej pamięci. Aby tego uniknąć, tekstury są wstępnie kompresowane i dynamicznie dekompresowane w razie potrzeby. A kompresja i dekompresja to obliczenia ciągłe. Aby pracować z teksturami, nVidia udostępniła pakiet NVIDIA Texture Tools. Obsługuje wydajną kompresję i dekompresję tekstur DirectX (tzw. format HFC). Druga wersja tego pakietu może pochwalić się obsługą algorytmów kompresji BC4 i BC5 zaimplementowanych w technologii DirectX 11. Ale najważniejsze jest to, że NVIDIA Texture Tools 2 obsługuje CUDA. Według nVidii daje to 12-krotny wzrost wydajności w zadaniach kompresji i dekompresji tekstur. A to oznacza, że ​​klatki rozgrywki będą ładowały się szybciej i zachwycą gracza swoim realizmem.

Pakiet NVIDIA Texture Tools 2 jest dostosowany do pracy z CUDA. Widoczny jest wzrost wydajności podczas kompresji i dekompresji tekstur.

Korzystanie z CUDA może znacznie poprawić wydajność monitoringu wideo.

Przetwarzanie strumienia wideo w czasie rzeczywistym

Mów, co chcesz, ale obecny świat, jeśli chodzi o szpiegostwo, jest znacznie bliższy światu Orwellowskiego Wielkiego Brata, niż się wydaje. Spojrzenie kamer wideo odczuwają zarówno kierowcy samochodów, jak i goście miejsc publicznych.

Pełno płynące rzeki informacji wideo wpływają do ośrodków jej przetwarzania i… wpadają w wąskie gardło – osobę. To on w większości przypadków jest ostatecznością, która monitoruje świat wideo. Co więcej, agencja nie jest najbardziej wydajna. Mruganie, rozkojarzenie i dążenie do zaśnięcia.

Dzięki CUDA możliwe stało się zaimplementowanie algorytmów do jednoczesnego śledzenia wielu obiektów w strumieniu wideo. W tym przypadku proces odbywa się w czasie rzeczywistym, a wideo jest pełne 30 kl./s. W porównaniu z implementacją takiego algorytmu na nowoczesnych procesorach wielordzeniowych, CUDA daje dwu-, trzykrotny wzrost wydajności, a to, jak widzisz, dużo.

Konwersja wideo, filtrowanie dźwięku

Badaboom Video Converter jest pierwszym, który używa CUDA do przyspieszenia konwersji.

Miło jest oglądać nową wypożyczalnię wideo w jakości FullHD i na dużym ekranie. Ale nie możesz zabrać ze sobą dużego ekranu w podróż, a kodek wideo FullHD zahaczy o energooszczędny procesor gadżetu mobilnego. Na ratunek przychodzi nawrócenie. Jednak większość z tych, którzy zetknęli się z nim w praktyce, narzeka na długi czas konwersji. Jest to zrozumiałe, proces jest rutynowy, odpowiedni do zrównoleglania, a jego wykonanie na procesorze nie jest zbyt optymalne.

Ale CUDA radzi sobie z tym z hukiem. Pierwszym znakiem jest konwerter Badaboom firmy Elevental. Deweloperzy Badaboom nie popełnili błędu, wybierając CUDA. Testy pokazują, że standardowe półtorej godziny filmu jest konwertowane do formatu iPhone/iPod Touch w mniej niż dwadzieścia minut. I to pomimo tego, że przy użyciu samego procesora proces ten trwa ponad godzinę.

Pomaga CUDA i profesjonalnym melomanom. Każdy z nich da połowę królestwa dla efektywnej zwrotnicy FIR - zestawu filtrów, które dzielą widmo dźwięku na kilka pasm. Proces ten jest bardzo czasochłonny i przy dużej ilości materiału audio sprawia, że ​​realizator dźwięku „zadymia” na kilka godzin. Wdrożenie zwrotnicy FIR opartej o CUDA przyspiesza jej pracę setki razy.

Przyszłość CUDA

Urzeczywistniając technologię GPGPU, CUDA nie zamierza spocząć na laurach. Jak wszędzie, w CUDA działa zasada odbicia: teraz nie tylko architektura procesorów wideo nVidia wpływa na rozwój wersji CUDA SDK, ale sama technologia CUDA zmusza nVidię do zrewidowania architektury jej chipów. Przykładem takiego odbicia jest platforma nVidia ION. Jego druga wersja jest specjalnie zoptymalizowana do rozwiązywania zadań CUDA. A to oznacza, że ​​nawet w stosunkowo niedrogich rozwiązaniach sprzętowych konsumenci otrzymają całą moc i genialne funkcje CUDA.

– zestaw niskopoziomowych interfejsów programistycznych ( API) do tworzenia gier i innych wydajnych aplikacji multimedialnych. Zawiera wsparcie dla wysokiej wydajności 2D- I 3D-grafika, dźwięk i urządzenia wejściowe.

Direct3D (D3D) – interfejs wyjściowy 3D prymitywne(ciała geometryczne). Zawarte w .

OpenGL(z angielskiego. Otwórz bibliotekę grafik, dosłownie - otwarta biblioteka graficzna) to specyfikacja, która definiuje niezależny od języka programowania, wieloplatformowy interfejs programowania do pisania aplikacji, które wykorzystują dwuwymiarową i trójwymiarową grafikę komputerową. Zawiera ponad 250 funkcji do rysowania złożonych scen 3D z prostych prymitywów. Znajduje zastosowanie w tworzeniu gier wideo, wirtualnej rzeczywistości, wizualizacji w badaniach naukowych. Na platformie Okna konkuruje z .

OpenCL(z angielskiego. Otwarty język komputerowy, dosłownie - otwarty język komputerowy) - struktura(szkielet systemu oprogramowania) do pisania programów komputerowych związanych z przetwarzaniem równoległym na różnych grafikach ( GPU) I ( ). Do ram OpenCL zawiera język programowania i interfejs programowania aplikacji ( API). OpenCL zapewnia równoległość na poziomie instrukcji i na poziomie danych i jest implementacją techniki GPU.

GPU(w skrócie z angielskiego. Jednostki przetwarzania grafiki ogólnego przeznaczenia, dosłownie - GPU ogólnego przeznaczenia) - technika używania procesora graficznego karty graficznej do ogólnych obliczeń, która jest zwykle przeprowadzana.

shader(Język angielski) shader) to program do konstruowania cieni na zsyntetyzowanych obrazach, używany w grafice trójwymiarowej do określania ostatecznych parametrów obiektu lub obrazu. Z reguły zawiera opis absorpcji i rozpraszania światła, mapowania tekstury, odbicia i załamania, cieniowania, przemieszczenia powierzchni i efektów post-processingu o dowolnej złożoności. Złożone powierzchnie można renderować za pomocą prostych kształtów geometrycznych.

wykonanie(Język angielski) wykonanie) - wizualizacja w grafice komputerowej procesu pozyskiwania obrazu z modelu za pomocą oprogramowania.

SDK(w skrócie z angielskiego. Zestaw programistyczny) to zestaw narzędzi programistycznych.

procesor(w skrócie z angielskiego. Jednostka centralna, dosłownie - centralne / główne / główne urządzenie obliczeniowe) - centralne (mikro);urządzenie wykonujące instrukcje maszynowe; sprzęt, który odpowiada za wykonywanie operacji obliczeniowych (nadanych przez system operacyjny i oprogramowanie aplikacyjne) oraz koordynację pracy wszystkich urządzeń.

GPU(w skrócie z angielskiego. Jednostka przetwarzania grafiki, dosłownie - graficzne urządzenie obliczeniowe) - procesor graficzny; oddzielne urządzenie lub konsola do gier, która wykonuje renderowanie grafiki (wizualizację). Nowoczesne procesory graficzne są bardzo wydajne w przetwarzaniu i realistycznym renderowaniu grafiki komputerowej. Procesor graficzny w nowoczesnych kartach wideo służy jako akcelerator grafiki 3D, ale w niektórych przypadkach może być również używany do obliczeń ( GPU).

Problemy procesor

Przez długi czas wzrost wydajności tradycyjnych wynikał głównie z konsekwentnego wzrostu częstotliwości taktowania (około 80% wydajności determinowała częstotliwość taktowania) przy jednoczesnym wzroście liczby tranzystorów na pojedynczym żeton. Jednak dalszy wzrost częstotliwości taktowania (przy częstotliwości taktowania powyżej 3,8 GHz chipy po prostu się przegrzewają!) opiera się wielu fundamentalnym barierom fizycznym (ponieważ proces technologiczny prawie zbliżył się do wielkości atomu: , a wielkość atomu krzemu wynosi około 0,543 nm):

Po pierwsze, wraz ze spadkiem wielkości kryształu i wzrostem częstotliwości zegara wzrasta prąd upływu tranzystorów. Prowadzi to do wzrostu zużycia energii i wzrostu emisji ciepła;

Po drugie, korzyści płynące z wyższych częstotliwości zegara są częściowo niwelowane przez opóźnienia dostępu do pamięci, ponieważ czasy dostępu do pamięci nie odpowiadają rosnącym częstotliwościom zegara;

Po trzecie, w przypadku niektórych aplikacji tradycyjne architektury szeregowe stają się nieefektywne w miarę wzrostu prędkości zegara z powodu tak zwanego „wąskiego gardła von Neumanna”, wąskiego gardła wydajności wynikającego z sekwencyjnego przepływu obliczeń. Jednocześnie zwiększają się opóźnienia transmisji sygnału rezystancyjno-pojemnościowego, co jest dodatkowym wąskim gardłem związanym ze wzrostem częstotliwości taktowania.

Rozwój GPU

Równolegle z rozwojem GPU:

…

Listopad 2008 - Intel wprowadził linię 4-rdzeniową Intel Core i7 w oparciu o mikroarchitekturę nowej generacji Nehalem. Procesory działają z częstotliwością zegara 2,6-3,2 GHz. Wykonany w technologii procesowej 45 nm.

Grudzień 2008 - Rozpoczęcie dostaw czterordzeniowych AMD Phenom II 940(kryptonim - Deneb). Pracuje na częstotliwości 3 GHz, produkowany jest w technologii procesowej 45 nm.

…

maj 2009 - firma AMD wprowadził wersję GPU ATI Radeon HD 4890 z taktowaniem rdzenia zwiększonym z 850 MHz do 1 GHz. To jest pierwszy graficzny procesor działający z częstotliwością 1 GHz. Moc obliczeniowa chipa, ze względu na wzrost częstotliwości, wzrosła z 1,36 do 1,6 teraflopów. Procesor zawiera 800 (!) rdzeni, obsługuje pamięć wideo GDDR5, DirectX 10.1, ATI CrossFireX i wszystkie inne technologie związane z nowoczesnymi modelami kart graficznych. Chip wykonany jest w oparciu o technologię 55 nm.

Główne różnice GPU

Cechy charakterystyczne GPU(w porównaniu z ) są:

– architekturę maksymalnie ukierunkowaną na zwiększenie szybkości obliczania tekstur i złożonych obiektów graficznych;

to szczytowa moc typowego GPU znacznie wyższy niż ;

– dzięki dedykowanej architekturze potoku, GPU znacznie wydajniejszy w przetwarzaniu informacji graficznych niż .

„Kryzys gatunku”

„Kryzys gatunku” dla dojrzał do 2005 r. - wtedy się pojawiły. Jednak pomimo rozwoju technologii, wzrost produktywności konwencjonalnej znacznie się zmniejszyła. W tym samym czasie wydajność GPU nadal się rozwija. Tak więc do 2003 roku ta rewolucyjna idea skrystalizowała się - wykorzystać moc obliczeniową grafiki. Procesory GPU stały się aktywnie wykorzystywane do obliczeń „niegraficznych” (symulacja fizyczna, przetwarzanie sygnałów, matematyka/geometria obliczeniowa, operacje na bazach danych, biologia obliczeniowa, ekonomia obliczeniowa, wizja komputerowa itp.).

Głównym problemem był brak standardowego interfejsu do programowania GPU. Wykorzystywane przez programistów OpenGL lub Direct3D ale to było bardzo wygodne. Korporacja NVIDIA(jeden z największych producentów procesorów graficznych, multimedialnych i komunikacyjnych, a także bezprzewodowych procesorów multimedialnych; założony w 1993 r.) był zaangażowany w opracowanie pewnego ujednoliconego i wygodnego standardu - i wprowadził tę technologię CUDA.

Jak to się zaczęło

2006 - NVIDIA demonstruje CUDA™; początek rewolucji w informatyce GPU.

2007 - NVIDIA wypuszcza architekturę CUDA(orginalna wersja CUDA SDK została zaprezentowana 15 lutego 2007); nominacja " Najlepsze nowe» z magazynu Popularna nauka oraz „Wybór czytelników” z publikacji HPCWire.

2008 - technologia NVIDIA CUDA wygrał w nominacji „Doskonałość techniczna” od Magazyn PC.

Co się stało CUDA

CUDA(w skrócie z angielskiego. Zunifikowana architektura urządzeń obliczeniowych, dosłownie - zunifikowana architektura obliczeniowa urządzeń) - architektura (zestaw oprogramowania i sprzętu), która pozwala produkować na GPU obliczenia ogólnego przeznaczenia GPU faktycznie działa jako potężny koprocesor.

Technologia NVIDIA CUDA™ jest jedynym środowiskiem programistycznym w języku programowania C, który umożliwia programistom tworzenie oprogramowania do rozwiązywania złożonych problemów obliczeniowych w krótszym czasie dzięki mocy obliczeniowej procesorów graficznych. Na świecie pracują już miliony GPU ze wsparciem CUDA, a tysiące programistów już korzysta (za darmo!) z narzędzi CUDA przyspieszać aplikacje i rozwiązywać najbardziej złożone zadania wymagające dużych zasobów — od kodowania wideo i audio po poszukiwanie ropy i gazu, modelowanie produktów, obrazowanie medyczne i badania naukowe.

CUDA daje programiście możliwość, według własnego uznania, organizowania dostępu do zestawu instrukcji akceleratora graficznego i zarządzania jego pamięcią, organizowania na nim złożonych obliczeń równoległych. Obsługiwany akcelerator graficzny CUDA staje się potężną programowalną otwartą architekturą, taką jak dzisiejsza. Wszystko to zapewnia deweloperowi niskopoziomowy, rozproszony i szybki dostęp do sprzętu, dzięki czemu CUDA niezbędną podstawę do budowy poważnych narzędzi wysokiego poziomu, takich jak kompilatory, debuggery, biblioteki matematyczne, platformy programowe.

Uralsky, główny specjalista ds. technologii NVIDIA, porównując GPU I , mówi tak: - To SUV. Podróżuje zawsze i wszędzie, ale niezbyt szybko. ALE GPU to samochód sportowy. Na złej drodze po prostu nigdzie nie pojedzie, ale da dobrą relację - i pokaże całą swoją prędkość, o której SUV nigdy nie marzył!..».

Możliwości technologiczne CUDA

Urządzenia do przekształcania komputerów osobistych w małe superkomputery znane są od dawna. W latach 80. ubiegłego wieku na rynku oferowano tzw. transputery, które umieszczano w powszechnych wówczas gniazdach rozszerzeń ISA. Początkowo ich wydajność w odpowiednich zadaniach była imponująca, ale potem wzrost wydajności procesorów uniwersalnych przyspieszył, umocniły swoją pozycję w obliczeniach równoległych, a transputery nie miały sensu. Chociaż takie urządzenia nadal istnieją, są to różnorodne wyspecjalizowane akceleratory. Często jednak zakres ich zastosowania jest wąski i takie akceleratory nie są powszechnie stosowane.

Ale ostatnio pałeczka obliczeń równoległych przeniosła się na rynek masowy, w taki czy inny sposób związany z grami trójwymiarowymi. Urządzenia ogólnego przeznaczenia z wielordzeniowymi procesorami do równoległego przetwarzania wektorowego używane w grafice 3D osiągają wysoką wydajność szczytową, której nie są w stanie osiągnąć procesory ogólnego przeznaczenia. Oczywiście maksymalna prędkość osiągana jest tylko w wielu wygodnych zadaniach i ma pewne ograniczenia, ale takie urządzenia już zaczęły być szeroko stosowane w obszarach, do których nie były pierwotnie przeznaczone. świetny przykład Takim procesorem równoległym jest procesor Cell opracowany przez sojusz Sony-Toshiba-IBM i używany w konsoli do gier Sony PlayStation 3, a także we wszystkich nowoczesnych kartach graficznych liderów rynku - Nvidii i AMD.

Dzisiaj Cella nie tkniemy, choć pojawił się wcześniej i jest uniwersalnym procesorem z dodatkowymi możliwościami wektorowymi, to dziś o nim nie mówimy. W przypadku akceleratorów wideo 3D pierwsze niegraficzne technologie obliczeniowe ogólnego przeznaczenia GPGPU (General-Purpose Computation na GPU) pojawiły się kilka lat temu. W końcu nowoczesne chipy wideo zawierają setki matematycznych jednostek wykonawczych, a tę moc można wykorzystać do znacznego przyspieszenia wielu wymagających obliczeniowo aplikacji. A obecne generacje procesorów graficznych mają wystarczająco elastyczną architekturę, która wraz z wysokopoziomowymi językami programowania i architekturami sprzętowo-programowymi, takimi jak ta omówiona w tym artykule, otwiera te możliwości i czyni je znacznie bardziej dostępnymi.

Stworzenie GPCPU było spowodowane pojawieniem się wystarczająco szybkich i elastycznych programów cieniujących, które są w stanie wykonywać nowoczesne chipy wideo. Twórcy postanowili, aby GPU obliczało nie tylko obraz w aplikacjach 3D, ale było również wykorzystywane w innych równoległych obliczeniach. W GPGPU wykorzystano do tego graficzne API: OpenGL i Direct3D, gdy dane były przesyłane do układu wideo w postaci tekstur, oraz programy obliczeniowe załadowane jako shadery. Wadami tej metody są stosunkowo duża złożoność programowania, niska prędkość wymiana danych między CPU i GPU oraz inne ograniczenia, które omówimy później.

Obliczenia na GPU ewoluowały i ewoluują bardzo szybko. Idąc dalej, dwóch głównych producentów układów wideo, Nvidia i AMD, opracowało i ogłosiło odpowiednie platformy o nazwach CUDA (Compute Unified Device Architecture) i CTM (Close To Metal lub AMD Stream Computing). W przeciwieństwie do poprzednich modeli programowania GPU, zostały one wykonane z bezpośrednim dostępem do możliwości sprzętowych kart graficznych. Platformy nie są ze sobą kompatybilne, CUDA jest rozszerzeniem języka programowania C, a CTM jest maszyna wirtualna, wykonując kod asemblera. Jednak obie platformy wyeliminowały niektóre z ważnych ograniczeń poprzednich modeli GPGPU przy użyciu tradycyjnego potoku graficznego i odpowiednich interfejsów Direct3D lub OpenGL.

Oczywiście otwarte standardy wykorzystujące OpenGL wydają się być najbardziej przenośne i uniwersalne, pozwalają na użycie tego samego kodu dla chipów wideo różnych producentów. Ale takie metody mają wiele wad, są znacznie mniej elastyczne i nie tak wygodne w użyciu. Ponadto uniemożliwiają korzystanie z określonych funkcji niektórych kart wideo, takich jak szybka współdzielona (współdzielona) pamięć obecna w nowoczesnych procesorach obliczeniowych.

Dlatego Nvidia wypuściła platformę CUDA, język programowania podobny do C, z własnym kompilatorem i bibliotekami do obliczeń GPU. Oczywiście napisanie optymalnego kodu dla chipów wideo wcale nie jest takie proste i to zadanie wymaga długiej ręcznej pracy, ale CUDA dopiero odsłania wszystkie możliwości i daje programiście większą kontrolę nad możliwościami sprzętowymi GPU. Ważne jest, aby obsługa Nvidia CUDA była dostępna dla układów G8x, G9x i GT2xx używanych w kartach graficznych Geforce 8, 9 i 200, które są bardzo rozpowszechnione. Została wydana ostateczna wersja CUDA 2.0, która ma kilka nowych funkcji, takich jak obsługa obliczeń podwójnej precyzji. CUDA jest dostępny w 32-bitowych i 64-bitowych systemach operacyjnych Linux, Windows i MacOS X.

Różnica między procesorem a procesorem graficznym w obliczeniach równoległych

Wzrost częstotliwości procesorów uniwersalnych napotkał fizyczne ograniczenia i wysoki pobór mocy, a ich wydajność coraz bardziej wzrasta dzięki umieszczeniu kilku rdzeni w jednym chipie. Sprzedawane obecnie procesory zawierają tylko do czterech rdzeni (dalszy wzrost nie będzie już szybki) i są przeznaczone do ogólnych zastosowań, wykorzystują MIMD - wielokrotny przepływ instrukcji i danych. Każdy rdzeń działa niezależnie od pozostałych, wykonując różne instrukcje dla różnych procesów.

Wyspecjalizowane możliwości wektorowe (SSE2 i SSE3) dla wektorów 4-składnikowych (zmiennoprzecinkowych o pojedynczej precyzji) i dwuskładnikowych (o podwójnej precyzji) pojawiły się w procesorach ogólnego przeznaczenia przede wszystkim ze względu na zwiększone wymagania aplikacji graficznych. Dlatego do niektórych zadań bardziej opłacalne jest wykorzystanie GPU, ponieważ zostały one pierwotnie dla nich stworzone.

Na przykład w chipach wideo Nvidii jednostka główna jest wieloprocesorem z ośmioma do dziesięciu rdzeniami i setkami jednostek ALU w sumie, kilkoma tysiącami rejestrów i niewielką ilością pamięci współdzielonej. Dodatkowo karta graficzna zawiera szybką pamięć globalną z dostępem do niej przez wszystkie wieloprocesory, pamięć lokalną w każdym wieloprocesorze oraz specjalną pamięć na stałe.

Co najważniejsze, te wieloprocesorowe rdzenie w GPU to rdzenie SIMD (pojedynczy strumień instrukcji, wiele strumieni danych). A te rdzenie wykonują te same instrukcje w tym samym czasie, ten styl programowania jest powszechny w algorytmach graficznych i wielu zadaniach naukowych, ale wymaga specyficznego programowania. Ale takie podejście pozwala na zwiększenie liczby jednostek wykonawczych dzięki ich uproszczeniu.

Wymieńmy więc główne różnice między architekturami CPU i GPU. Rdzenie procesora są zaprojektowane do wykonywania pojedynczego strumienia sekwencyjnych instrukcji z maksymalną wydajnością, podczas gdy procesory graficzne są zaprojektowane do szybkiego wykonywania dużej liczby równoległych strumieni instrukcji. Procesory ogólnego przeznaczenia są zoptymalizowane pod kątem osiągnięcia wysokiej wydajności na pojedynczym strumieniu instrukcji, który przetwarza zarówno liczby całkowite, jak i zmiennoprzecinkowe. Dostęp do pamięci jest losowy.

Projektanci procesorów starają się uzyskać jak najwięcej instrukcji wykonywanych równolegle w celu zwiększenia wydajności. W tym celu, począwszy od procesorów Intel Pentium, pojawiło się wykonywanie superskalarne, zapewniające wykonanie dwóch instrukcji na zegar, a Pentium Pro wyróżniał się wykonywaniem instrukcji poza kolejnością. Jednak równoległe wykonywanie sekwencyjnego strumienia instrukcji ma pewne podstawowe ograniczenia, a poprzez zwiększenie liczby jednostek wykonawczych nie można osiągnąć wielokrotnego zwiększenia szybkości.

Chipy wideo od samego początku działają w prosty i równoległy sposób. Układ wideo pobiera grupę wielokątów na wejściu, wykonuje wszystkie niezbędne operacje i generuje piksele na wyjściu. Przetwarzanie wielokątów i pikseli jest niezależne, mogą być przetwarzane równolegle, niezależnie od siebie. Dlatego też, ze względu na z natury równoległą organizację pracy w GPU, stosuje się dużą liczbę jednostek wykonawczych, które są łatwe do załadowania, w przeciwieństwie do sekwencyjnego przepływu instrukcji dla CPU. Ponadto nowoczesne procesory graficzne mogą wykonywać więcej niż jedną instrukcję na zegar (problem podwójny). Tak więc architektura Tesli, pod pewnymi warunkami, uruchamia jednocześnie operacje MAD+MUL lub MAD+SFU.

GPU różni się od CPU także zasadami dostępu do pamięci. W GPU jest podłączona i łatwo przewidywalna - jeśli z pamięci zostanie odczytany teksel tekstur, to po chwili przyjdzie czas na sąsiednie teksele. Tak, a przy nagrywaniu to samo – piksel jest zapisywany do bufora ramki, a po kilku cyklach zapisywany będzie ten znajdujący się obok. Dlatego organizacja pamięci różni się od tej używanej w procesorze. A chip wideo, w przeciwieństwie do procesorów uniwersalnych, po prostu nie potrzebuje dużej pamięci podręcznej, a tekstury wymagają zaledwie kilku (do 128-256 w obecnych GPU) kilobajtów.

A sama praca z pamięcią dla GPU i CPU jest nieco inna. Tak więc nie wszystkie procesory mają wbudowane kontrolery pamięci, a wszystkie procesory graficzne mają zwykle kilka kontrolerów, do ośmiu 64-bitowych kanałów w układzie Nvidia GT200. Ponadto karty graficzne wykorzystują szybszą pamięć, w wyniku czego chipy wideo mają wielokrotnie większą dostępną przepustowość pamięci, co jest również bardzo ważne w przypadku obliczeń równoległych, które operują ogromnymi strumieniami danych.

W procesorach ogólnego przeznaczenia duża liczba tranzystorów i obszarów chipów trafia do buforów instrukcji, przewidywania gałęzi sprzętowych i ogromnych ilości pamięci podręcznej na chipie. Wszystkie te bloki sprzętowe są potrzebne, aby przyspieszyć wykonanie kilku strumieni instrukcji. Chipy wideo zużywają tranzystory na macierze jednostek wykonawczych, jednostki sterujące przepływem, małą pamięć współdzieloną i wielokanałowe kontrolery pamięci. Powyższe nie przyspiesza wykonywania poszczególnych wątków, pozwala chipowi obsłużyć kilka tysięcy wątków jednocześnie wykonywanych na chipie i wymagających dużej przepustowości pamięci.

O różnicach w buforowaniu. Procesory ogólnego przeznaczenia używają pamięci podręcznej, aby zwiększyć wydajność, zmniejszając opóźnienia dostępu do pamięci, podczas gdy procesory GPU używają pamięci podręcznej lub pamięci współdzielonej, aby zwiększyć przepustowość. Procesory zmniejszają opóźnienia w dostępie do pamięci dzięki dużym pamięciom podręcznym i przewidywaniu gałęzi kodu. Te elementy sprzętowe zajmują większość obszaru chipa i zużywają dużo energii. Układy wideo omijają problem opóźnień dostępu do pamięci, wykonując jednocześnie tysiące wątków — podczas gdy jeden z wątków czeka na dane z pamięci, układ wideo może wykonywać obliczenia innego wątku bez czekania i opóźnień.

Istnieje również wiele różnic w obsłudze wielowątkowości. Procesor wykonuje 1-2 wątki obliczeniowe na rdzeń procesora, a chipy wideo mogą obsługiwać do 1024 wątków na wieloprocesor, z których kilka jest w chipie. A jeśli przełączanie procesora z jednego wątku na inny kosztuje setki cykli, to GPU przełącza kilka wątków w jednym cyklu.

Ponadto procesory wykorzystują bloki SIMD (pojedyncza instrukcja, wiele danych) do obliczeń wektorowych, a procesory graficzne używają SIMT (pojedyncza instrukcja, wiele wątków) do przetwarzania wątków skalarnych. SIMT nie wymaga od programisty konwersji danych na wektory i umożliwia dowolne rozgałęzianie w strumienie.

Krótko mówiąc, możemy powiedzieć, że w przeciwieństwie do nowoczesnych uniwersalnych procesorów, chipy wideo są przeznaczone do obliczeń równoległych z dużą liczbą operacji arytmetycznych. A znacznie większa liczba tranzystorów GPU pracuje zgodnie z ich przeznaczeniem – przetwarzaniem tablic danych, a nie steruje wykonywaniem (sterowaniem przepływem) kilku sekwencyjnych wątków obliczeniowych. To jest wykres pokazujący, ile miejsca w CPU i GPU zajmuje różnorodna logika:

W efekcie podstawą efektywnego wykorzystania mocy GPU w obliczeniach naukowych i innych niegraficznych jest zrównoleglanie algorytmów na setki jednostek wykonawczych dostępnych w chipach wideo. Na przykład wiele zastosowań modelowania molekularnego dobrze nadaje się do obliczeń na chipach wideo, wymagają one dużej mocy obliczeniowej i dlatego są wygodne do obliczeń równoległych. A użycie wielu procesorów graficznych zapewnia jeszcze większą moc obliczeniową do rozwiązywania takich problemów.

Wykonywanie obliczeń na GPU daje doskonałe wyniki w algorytmach wykorzystujących równoległe przetwarzanie danych. Oznacza to, że ta sama sekwencja operacji matematycznych jest stosowana do dużej ilości danych. W takim przypadku najlepsze wyniki uzyskuje się, gdy stosunek liczby instrukcji arytmetycznych do liczby dostępów do pamięci jest wystarczająco duży. Stwarza to mniejsze wymagania w zakresie kontroli przepływu, a duża gęstość matematyki i duża ilość danych eliminuje potrzebę stosowania dużych pamięci podręcznych, jak w przypadku procesora.

W wyniku wszystkich opisanych powyżej różnic teoretyczna wydajność układów wideo znacznie przewyższa wydajność procesora. Nvidia przedstawia następujący wykres wzrostu wydajności procesora i karty graficznej w ciągu ostatnich kilku lat:

Oczywiście dane te nie są pozbawione chytrości. Rzeczywiście, na CPU o wiele łatwiej jest osiągnąć teoretyczne wartości w praktyce, a liczby podane są dla pojedynczej precyzji w przypadku GPU i podwójnej precyzji w przypadku CPU. W każdym razie pojedyncza precyzja wystarczy do niektórych równoległych zadań, a różnica w szybkości między procesorami uniwersalnymi i graficznymi jest bardzo duża, dlatego gra jest warta świeczki.

Pierwsze próby zastosowania obliczeń na GPU

Chipy wideo były używane od dawna w równoległych obliczeniach matematycznych. Pierwsze próby stworzenia takiej aplikacji były niezwykle prymitywne i ograniczały się do wykorzystania niektórych funkcji sprzętowych, takich jak rasteryzacja i buforowanie Z. Ale w obecnym stuleciu, wraz z pojawieniem się shaderów, zaczęli przyspieszać obliczanie macierzy. W 2003 roku w SIGGRAPH wydzielono osobną sekcję dla obliczeń GPU i nazwano ją GPGPU (General-Purpose Computation na GPU) – uniwersalne przetwarzanie GPU.

Najbardziej znanym BrookGPU jest kompilator języka programowania Brook, zaprojektowany do wykonywania obliczeń niegraficznych na GPU. Przed jego pojawieniem się programiści wykorzystujący do obliczeń możliwości chipów wideo wybrali jeden z dwóch popularnych API: Direct3D lub OpenGL. To poważnie ograniczyło użycie GPU, ponieważ grafika 3D wykorzystuje shadery i tekstury, o których programiści równolegli nie muszą wiedzieć, używają wątków i rdzeni. Brook był w stanie ułatwić im zadanie. Te rozszerzenia przesyłania strumieniowego do języka C, opracowane na Uniwersytecie Stanforda, ukryły interfejs API 3D przed programistami i zaprezentowały układ wideo jako równoległy koprocesor. Kompilator przeanalizował plik .br z kodem C++ i rozszerzeniami, tworząc kod połączony z biblioteką obsługującą DirectX, OpenGL lub x86.

Oczywiście Brook miał wiele niedociągnięć, nad którymi się zastanawiamy i które omówimy bardziej szczegółowo później. Ale już sam jego wygląd spowodował znaczny wzrost zainteresowania tej samej Nvidii i ATI inicjatywą obliczeń GPU, ponieważ rozwój tych możliwości poważnie zmienił rynek w przyszłości, otwierając zupełnie nowy jego sektor - obliczenia równoległe oparte na chipy wideo.

Co więcej, niektórzy badacze z projektu Brook dołączyli do zespołu programistów Nvidii, aby wprowadzić strategię równoległego przetwarzania sprzętowo-programowego, otwierając nowy udział w rynku. Główną zaletą tej inicjatywy Nvidii było to, że programiści doskonale znają wszystkie możliwości swoich procesorów graficznych w najdrobniejszych szczegółach i nie ma potrzeby korzystania z graficznego API, a ze sprzętem można pracować bezpośrednio za pomocą sterownika. Efektem wysiłków tego zespołu jest Nvidia CUDA (Compute Unified Device Architecture), nowa architektura sprzętowo-programowa do obliczeń równoległych na GPU Nvidia, która jest przedmiotem tego artykułu.

Obszary zastosowania obliczeń równoległych na GPU

Aby zrozumieć, jakie korzyści przynosi przeniesienie obliczeń na chipy wideo, przedstawimy średnie dane uzyskane przez badaczy na całym świecie. Średnio przy przenoszeniu obliczeń na GPU w wielu zadaniach przyspieszenie osiąga się 5-30 razy w porównaniu do szybkich procesorów uniwersalnych. Największe liczby (rzędu przyspieszenia 100x, a nawet więcej!) są osiągane na kodzie, który nie jest zbyt dobrze przystosowany do obliczeń z użyciem bloków SSE, ale jest całkiem wygodny dla GPU.

To tylko kilka przykładów przyspieszenia kodu syntetycznego na GPU w porównaniu z kodem wektorowym SSE na CPU (według Nvidii):

• Mikroskopia fluorescencyjna: 12x;
• Dynamika molekularna (obliczona siła niezwiązana): 8-16x;
• Elektrostatyka (bezpośrednie i wielopoziomowe sumowanie kulombowskie): 40-120x i 7x.

A jest to płyta, którą Nvidia bardzo kocha, pokazując ją na wszystkich prezentacjach, o których bardziej szczegółowo przyjrzymy się w drugiej części artykułu, poświęconej konkretnym przykładom praktycznych zastosowań obliczeń CUDA:

Jak widać, liczby są bardzo atrakcyjne, zwłaszcza 100-150-krotne zyski robią wrażenie. W następnym artykule CUDA przyjrzymy się bliżej niektórym z tych liczb. A teraz wymieniamy główne zastosowania, w których obecnie wykorzystywane są obliczenia na GPU: analiza i przetwarzanie obrazów i sygnałów, symulacja fizyczna, matematyka obliczeniowa, biologia obliczeniowa, obliczenia finansowe, bazy danych, dynamika gazów i cieczy, kryptografia, radioterapia adaptacyjna, astronomia , przetwarzanie dźwięku, bioinformatyka, symulacje biologiczne, wizja komputerowa, eksploracja danych, kino cyfrowe i telewizja, symulacje elektromagnetyczne, systemy informacji geograficznej, zastosowania wojskowe, planowanie górnicze, dynamika molekularna, rezonans magnetyczny (MRI), sieci neuronowe, badania oceanograficzne, cząstki fizyka, symulacja składania białek, chemia kwantowa, śledzenie promieni, obrazowanie, radar, symulacja zbiorników, sztuczna inteligencja, analiza danych satelitarnych, eksploracja sejsmiczna, chirurgia, ultradźwięki, wideokonferencje.

Szczegóły wielu aplikacji można znaleźć na stronie Nvidii w dziale . Jak widać, lista jest całkiem spora, ale to nie wszystko! Można to kontynuować i z pewnością możemy założyć, że w przyszłości pojawią się inne obszary zastosowań obliczeń równoległych na chipach wideo, o których wciąż nie mamy pojęcia.

Możliwości NVIDIA CUDA

Technologia CUDA to programowa i sprzętowa architektura obliczeniowa Nvidii oparta na rozszerzeniu języka C, która umożliwia dostęp do zestawu instrukcji akceleratora graficznego i zarządzanie jego pamięcią w obliczeniach równoległych. CUDA pomaga zaimplementować algorytmy, które można zaimplementować na procesorach graficznych akceleratorów wideo Geforce ósmej generacji i starszych (serie Geforce 8, Geforce 9, Geforce 200), a także Quadro i Tesla.

Chociaż złożoność programowania GPU z CUDA jest dość duża, jest ona niższa niż we wczesnych rozwiązaniach GPGPU. Takie programy wymagają partycjonowania aplikacji na wiele procesorów wieloprocesorowych, podobnie jak w przypadku programowania MPI, ale bez współdzielenia danych przechowywanych we współdzielonej pamięci wideo. A ponieważ programowanie CUDA dla każdego wieloprocesora jest podobne do programowania OpenMP, wymaga dobrego zrozumienia organizacji pamięci. Ale oczywiście złożoność tworzenia i przenoszenia do CUDA w dużym stopniu zależy od aplikacji.

Zestaw deweloperski zawiera wiele przykładów kodu i jest dobrze udokumentowany. Proces uczenia się zajmie około dwóch do czterech tygodni osobom już zaznajomionym z OpenMP i MPI. API jest oparte na rozszerzonym języku C, a do tłumaczenia kodu z tego języka, CUDA SDK zawiera kompilator wiersza poleceń nvcc, oparty na otwartym kompilatorze Open64.

Podajemy główne cechy CUDA:

• zunifikowane rozwiązanie programowe i sprzętowe do obliczeń równoległych na chipach wideo Nvidia;
• szeroka gama obsługiwanych rozwiązań, od mobilnych po wielochipowe
• standardowy język programowania C;
• standardowe biblioteki do analizy numerycznej FFT (Fast Fourier Transform) i BLAS (Linear Algebra);
• zoptymalizowana wymiana danych między CPU i GPU;
• interakcja z graficznym API OpenGL i DirectX;
• obsługa 32- i 64-bitowych systemów operacyjnych: Windows XP, Windows Vista, Linux i MacOS X;
• zdolność do rozwoju na niskim poziomie.

Odnośnie obsługi systemów operacyjnych należy dodać, że wszystkie główne Dystrybucje Linuksa(Red Hat Enterprise Linux 3.x/4.x/5.x, SUSE Linux 10.x), ale według entuzjastów CUDA działa dobrze na innych kompilacjach: Fedora Core, Ubuntu, Gentoo itp.

Środowisko programistyczne CUDA (CUDA Toolkit) obejmuje:

• kompilator nvcc;
• biblioteki FFT i BLAS;
• profiler;
• debugger gdb dla GPU;
• Sterownik wykonawczy CUDA dołączony do standardowych sterowników Nvidii
• instrukcja programowania;
• CUDA Developer SDK (kod źródłowy, narzędzia i dokumentacja).

W przykładach kod źródłowy: równoległe sortowanie bitonic (bitonic sort), transpozycja macierzy, równoległe sumowanie prefiksów dużych tablic, splot obrazu, dyskretna transformata falkowa, przykład interakcji z OpenGL i Direct3D, wykorzystanie bibliotek CUBLAS i CUFFT, kalkulacja ceny opcji (formuła Blacka-Scholesa, model dwumianowy, metoda Monte Carlo), równoległy generator liczb losowych Mersenne Twister, obliczanie histogramu dużej tablicy, redukcja szumów, filtr Sobela (odnajdywanie krawędzi).

Korzyści i ograniczenia CUDA

Z punktu widzenia programisty potok graficzny to zestaw etapów przetwarzania. Blok geometrii generuje trójkąty, a blok rasteryzacji generuje piksele, które są wyświetlane na monitorze. Tradycyjny model programowania GPGPU wygląda następująco:

Aby przenieść obliczenia na GPU w ramach takiego modelu, potrzebne jest specjalne podejście. Nawet dodanie dwóch wektorów element po elemencie będzie wymagało narysowania kształtu na ekranie lub w buforze poza ekranem. Rysunek jest rasteryzowany, kolor każdego piksela jest obliczany zgodnie z danym programem (pixel shader). Program odczytuje dane wejściowe z tekstur dla każdego piksela, dodaje je i zapisuje w buforze wyjściowym. A wszystkie te liczne operacje są potrzebne do tego, co jest napisane w jednym operatorze w konwencjonalnym języku programowania!

Dlatego wykorzystanie GPGPU do obliczeń ogólnego przeznaczenia ma ograniczenie w postaci zbyt dużej złożoności, aby programiści mogli się uczyć. I jest wystarczająco dużo innych ograniczeń, ponieważ pixel shader to tylko formuła na zależność ostatecznego koloru piksela od jego współrzędnych, a pixel shader to język pisania tych formuł ze składnią podobną do C. Wczesne metody GPGPU to sprytna sztuczka, która pozwala wykorzystać moc GPU, ale bez żadnej wygody. Tam dane są reprezentowane przez obrazy (tekstury), a algorytm jest reprezentowany przez proces rasteryzacji. Należy zwrócić uwagę na bardzo specyficzny model pamięci i wykonania.

Architektura sprzętowa i programowa Nvidii do obliczeń na procesorach graficznych firmy Nvidia różni się od poprzednich modeli GPGPU tym, że umożliwia pisanie programów dla procesorów graficznych w prawdziwym C ze standardową składnią, wskaźnikami i koniecznością posiadania minimum rozszerzeń, aby uzyskać dostęp do zasobów obliczeniowych układów wideo. CUDA nie zależy od graficznych interfejsów API i ma pewne funkcje zaprojektowane specjalnie do przetwarzania ogólnego przeznaczenia.

Przewaga CUDA nad tradycyjnym podejściem do obliczeń GPGPU:

• interfejs programowania aplikacji CUDA oparty jest na standardowym języku programowania C z rozszerzeniami, co upraszcza proces uczenia się i implementacji architektury CUDA;
• CUDA zapewnia dostęp do 16 KB pamięci współdzielonej na wieloprocesor, która może być wykorzystana do zorganizowania pamięci podręcznej o większej przepustowości niż pobieranie tekstur;
• wydajniejszy transfer danych między systemem a pamięcią wideo
• nie ma potrzeby korzystania z graficznych interfejsów API z nadmiarowością i narzutem;
• liniowe adresowanie pamięci oraz gromadzenie i rozpraszanie, umiejętność pisania na dowolne adresy;
• wsparcie sprzętowe dla operacji na liczbach całkowitych i bitach.

Główne ograniczenia CUDA:

• brak obsługi rekurencji dla funkcji wykonywalnych;
• minimalna szerokość bloku to 32 wątki;
• autorska architektura CUDA należąca do Nvidii.

Słabości programowania z poprzednimi metodami GPGPU polegają na tym, że metody te nie używają jednostek wykonawczych Vertex Shader w poprzednich architekturach niezunifikowanych, dane są przechowywane w teksturach i wyprowadzane do bufora poza ekranem, a algorytmy wieloprzebiegowe używają jednostek cieniowania pikseli. Ograniczenia GPGPU obejmują: niewystarczająco wydajne wykorzystanie możliwości sprzętowych, ograniczenia przepustowości pamięci, brak operacji rozproszonych (tylko gromadzenie), obowiązkowe użycie graficznego API.

Główne zalety CUDA nad poprzednimi metodami GPGPU wynikają z faktu, że architektura ta została zaprojektowana tak, aby efektywnie wykorzystywać na GPU obliczenia inne niż graficzne i wykorzystuje język programowania C bez konieczności przenoszenia algorytmów do postaci wygodnej dla koncepcji grafiki rurociąg. CUDA oferuje nową ścieżkę obliczeniową GPU, która nie wykorzystuje graficznych interfejsów API, oferując losowy dostęp do pamięci (scatter lub collect). Taka architektura jest pozbawiona wad GPGPU i wykorzystuje wszystkie jednostki wykonawcze, a także rozszerza możliwości o matematykę liczb całkowitych i operacje przesunięcia bitowego.

Ponadto CUDA otwiera niektóre funkcje sprzętowe niedostępne w graficznych interfejsach API, takie jak pamięć współdzielona. Jest to niewielka ilość pamięci (16 kilobajtów na wieloprocesor), do której mają dostęp bloki wątków. Umożliwia buforowanie najczęściej używanych danych i może zapewnić wyższą wydajność niż przy użyciu pobierania tekstur do tego zadania. To z kolei zmniejsza czułość przepustowości algorytmów równoległych w wielu aplikacjach. Na przykład przydaje się do algebry liniowej, szybkiej transformacji Fouriera i filtrów przetwarzania obrazu.

Wygodniejszy w dostępie do CUDA i pamięci. Kod w graficznym API wyprowadza dane jako 32 wartości zmiennoprzecinkowe o pojedynczej precyzji (wartości RGBA jednocześnie do ośmiu celów renderowania) w predefiniowanych obszarach, a CUDA obsługuje zapis rozproszony - nieograniczoną liczbę rekordów pod dowolnym adresem. Takie zalety umożliwiają wykonanie niektórych algorytmów na GPU, których nie da się wydajnie zaimplementować metodami GPGPU opartymi na graficznym API.

Ponadto graficzne API muszą przechowywać dane w teksturach, co wymaga wcześniejszego upakowania dużych tablic w tekstury, co komplikuje algorytm i wymusza zastosowanie specjalnego adresowania. A CUDA pozwala na odczyt danych pod dowolnym adresem. Kolejną zaletą CUDA jest zoptymalizowana komunikacja między CPU i GPU. A dla programistów, którzy chcą uzyskać dostęp do niskiego poziomu (na przykład podczas pisania innego języka programowania), CUDA oferuje możliwość programowania w języku asemblera niskiego poziomu.

Historia rozwoju CUDA

Rozwój CUDA został ogłoszony wraz z układem G80 w listopadzie 2006 r. i wydaniem publiczna wersja beta CUDA SDK odbyło się w lutym 2007 roku. Wersja 1.0 została wydana w czerwcu 2007 roku, aby wprowadzić rozwiązania Tesla oparte na chipie G80 dla rynku obliczeń o wysokiej wydajności. Następnie pod koniec roku ukazała się wersja beta CUDA 1.1, która pomimo niewielkiego wzrostu numeru wersji wprowadziła sporo nowości.

Z tego, co pojawiło się w CUDA 1.1, możemy zauważyć włączenie funkcjonalności CUDA do zwykłych sterowników wideo Nvidii. Oznaczało to, że w wymaganiach dla dowolnego programu CUDA wystarczyło podać kartę graficzną Geforce serii 8 i wyższą oraz minimalną wersję sterownika 169.xx. Jest to bardzo ważne dla programistów, jeśli te warunki zostaną spełnione, programy CUDA będą działać dla każdego użytkownika. Dodano również wykonywanie asynchroniczne wraz z kopiowaniem danych (tylko dla układów G84, G86, G92 i wyższych), asynchroniczny transfer danych do pamięci wideo, operacje dostępu do pamięci atomowej, wsparcie dla 64-bitowych wersji Windows oraz możliwość multi -chip praca CUDA w trybie SLI.

Obecnie wersją dla rozwiązań opartych na GT200 jest CUDA 2.0, która została wydana wraz z linią Geforce GTX 200. Wersja beta została wydana wiosną 2008 roku. Druga wersja ma: obsługę obliczeń podwójnej precyzji (obsługa sprzętowa tylko dla GT200), wreszcie obsługiwane są Windows Vista (wersje 32 i 64-bitowe) i Mac OS X, dodano narzędzia do debugowania i profilowania, obsługiwane są tekstury 3D, zoptymalizowane transfer danych.

Jeśli chodzi o obliczenia z podwójną precyzją, ich szybkość na obecnej generacji sprzętu jest kilkakrotnie mniejsza niż pojedyncza precyzja. Powody są omówione w naszym. Implementacja tej obsługi w GT200 polega na tym, że bloki FP32 nie są wykorzystywane do uzyskiwania wyników w czterokrotnie wolniejszym tempie, do obsługi obliczeń FP64 Nvidia zdecydowała się na wykonanie dedykowanych bloków obliczeniowych. A w GT200 jest ich dziesięć razy mniej niż bloków FP32 (jeden blok podwójnej precyzji na każdy multiprocesor).

W rzeczywistości wydajność może być jeszcze niższa, ponieważ architektura jest zoptymalizowana pod kątem 32-bitowego odczytu z pamięci i rejestrów, ponadto podwójna precyzja nie jest potrzebna w aplikacjach graficznych, a w GT200 jest bardziej prawdopodobna. Tak, a współczesne czterordzeniowe procesory wykazują niewiele mniejszą realną wydajność. Ale będąc nawet 10 razy wolniejszym niż pojedyncza precyzja, ta obsługa jest przydatna w przypadku obwodów o mieszanej precyzji. Jedną z powszechnych technik jest uzyskiwanie początkowo przybliżonych wyników z pojedynczą precyzją, a następnie poprawianie ich z podwójną precyzją. Teraz można to zrobić bezpośrednio na karcie graficznej, bez wysyłania danych pośrednich do procesora.

Co dziwne, kolejna użyteczna funkcja CUDA 2.0 nie ma nic wspólnego z GPU. Dopiero teraz można skompilować kod CUDA w wysoce wydajny wielowątkowy kod SSE w celu szybkiego wykonania na procesorze. Oznacza to, że teraz ta funkcja nadaje się nie tylko do debugowania, ale także prawdziwy użytek w systemach bez karty graficznej Nvidia. W końcu wykorzystanie CUDA w normalnym kodzie jest ograniczone tym, że karty graficzne Nvidii, choć najbardziej popularne wśród dedykowanych rozwiązań wideo, nie są dostępne we wszystkich systemach. A przed wersją 2.0 w takich przypadkach należałoby napisać dwa różne kody: dla CUDA i osobno dla CPU. A teraz możesz uruchomić dowolny program CUDA na procesorze z wysoką wydajnością, aczkolwiek z mniejszą prędkością niż na chipach wideo.

Obsługiwane rozwiązania Nvidia CUDA

Wszystkie karty graficzne z obsługą CUDA mogą przyspieszyć najbardziej wymagające zadania, od przetwarzania audio i wideo po badania medyczne i naukowe. Jedynym prawdziwym ograniczeniem jest to, że wiele programów CUDA wymaga co najmniej 256 megabajtów pamięci wideo i jest to jedna z najważniejszych specyfikacji aplikacji CUDA.

Aktualną listę produktów obsługujących CUDA można znaleźć pod adresem . W momencie pisania tego tekstu obliczenia CUDA wspierały wszystkie produkty z serii Geforce 200, Geforce 9 i Geforce 8, w tym produkty mobilne, począwszy od Geforce 8400M, a także chipsety Geforce 8100, 8200 i 8300. Nowoczesne Quadro i wszystkie Tesla: S1070, C1060, C870, D870 i S870.

Szczególnie zwracamy uwagę, że wraz z nowymi kartami graficznymi Geforce GTX 260 i 280 ogłoszono odpowiednie rozwiązania obliczeniowe o wysokiej wydajności: Tesla C1060 i S1070 (pokazane na powyższym zdjęciu), które będą dostępne w sprzedaży jesienią tego roku. Zastosowano w nich ten sam GPU - GT200, w C1060 jeden, w S1070 - cztery. Ale w przeciwieństwie do rozwiązań do gier wykorzystują cztery gigabajty pamięci na chip. Z minusów być może niższa częstotliwość pamięci i przepustowość pamięci niż w przypadku kart do gier, zapewniająca 102 gigabajty / s na chip.

Skład Nvidii CUDA

CUDA zawiera dwa API: wysokiego poziomu (CUDA Runtime API) i niskiego poziomu (CUDA Driver API), chociaż nie można używać obu jednocześnie w jednym programie, należy użyć jednego lub drugiego. Wysokopoziomowy działa „na górze” niskopoziomowego, wszystkie wywołania uruchomieniowe są tłumaczone na proste instrukcje przetwarzane przez niskopoziomowe API Drivera. Ale nawet API „wysokiego poziomu” zakłada wiedzę na temat budowy i działania chipów wideo Nvidii, nie ma tam zbyt wysokiego poziomu abstrakcji.

Jest jeszcze jeden poziom, jeszcze wyższy - dwie biblioteki:

KUBLAS- wersja CUDA BLAS (Basic Linear Algebra Subprograms), przeznaczona do obliczania problemów algebry liniowej i korzystania z bezpośredniego dostępu do zasobów GPU;

MANKIET- Wersja CUDA biblioteki szybkiej transformacji Fouriera do obliczania szybkiej transformacji Fouriera, szeroko stosowanej w przetwarzaniu sygnałów. Obsługiwane są następujące typy przekształceń: złożone-złożone (C2C), rzeczywiste-złożone (R2C) i złożone-rzeczywiste (C2R).

Przyjrzyjmy się bliżej tym bibliotekom. CUBLAS to standardowe algorytmy algebry liniowej przetłumaczone na język CUDA, w chwili obecnej obsługiwany jest tylko pewien zestaw podstawowych funkcji CUBLAS. Biblioteka jest bardzo łatwa w użyciu: należy utworzyć macierz i obiekty wektorowe w pamięci wideo, wypełnić je danymi, wywołać wymagane funkcje CUBLAS i załadować wyniki z pamięci wideo z powrotem do pamięci systemowej. CUBLAS zawiera specjalne funkcje do tworzenia i niszczenia obiektów w pamięci GPU, a także do odczytywania i zapisywania danych do tej pamięci. Obsługiwane funkcje BLAS: poziomy 1, 2 i 3 dla liczb rzeczywistych, poziom 1 CGEMM dla złożonych. Poziom 1 to operacje wektorowo-wektorowe, poziom 2 to operacje na macierzach wektorowych, poziom 3 to operacje na macierzy.

CUFFT - wariant CUDA szybkiej transformacji Fouriera - szeroko stosowany i bardzo ważny w analizie sygnałów, filtrowaniu itp. CUFFT zapewnia prosty interfejs do wydajnych obliczeń FFT na procesorach graficznych Nvidia bez konieczności opracowywania niestandardowych FFT dla GPU. Wariant CUDA FFT obsługuje transformacje 1D, 2D i 3D złożonych i rzeczywistych danych, wykonywanie wsadowe dla wielu transformacji 1D równolegle, rozmiary transformacji 2D i 3D mogą mieścić się w zakresie , dla 1D obsługiwany jest rozmiar do 8 milionów elementów.

Podstawy tworzenia programów na CUDA

Aby zrozumieć poniższy tekst, powinieneś zrozumieć podstawowe cechy architektury chipów wideo Nvidii. GPU składa się z kilku klastrów jednostek tekstur (klaster przetwarzania tekstur). Każdy klaster składa się z powiększonego bloku pobieranych tekstur i dwóch lub trzech multiprocesorów strumieniowych, z których każdy składa się z ośmiu urządzenia komputerowe oraz dwa superfunkcjonalne bloki. Wszystkie instrukcje są wykonywane zgodnie z zasadą SIMD, gdy jedna instrukcja jest stosowana do wszystkich wątków w osnowie (termin z branży tekstylnej, w CUDA jest to grupa 32 wątków - minimalna ilość danych przetwarzanych przez wieloprocesory). Ta metoda wykonania została nazwana SIMT (pojedyncza instrukcja wiele wątków - jedna instrukcja i wiele wątków).

Każdy z multiprocesorów ma określone zasoby. Istnieje więc specjalna pamięć współdzielona o pojemności 16 kilobajtów na wieloprocesor. Ale to nie jest pamięć podręczna, ponieważ programista może jej użyć do dowolnych potrzeb, podobnie jak Sklep lokalny w SPU procesorów Cell. Ta pamięć współdzielona umożliwia wymianę informacji między wątkami tego samego bloku. Ważne jest, aby wszystkie wątki jednego bloku były zawsze wykonywane przez ten sam multiprocesor. A wątki z różnych bloków nie mogą wymieniać danych i musisz pamiętać o tym ograniczeniu. Pamięć współdzielona jest często przydatna, z wyjątkiem sytuacji, gdy wiele wątków uzyskuje dostęp do tego samego banku pamięci. Multiprocesory mogą również uzyskiwać dostęp do pamięci wideo, ale z większym opóźnieniem i gorszą przepustowością. Aby przyspieszyć dostęp i zmniejszyć częstotliwość uzyskiwania dostępu do pamięci wideo, wieloprocesory mają 8 kilobajtów pamięci podręcznej na stałe i dane tekstur.

Multiprocesor wykorzystuje 8192-16384 (odpowiednio dla G8x/G9x i GT2xx) rejestry wspólne dla wszystkich wątków wszystkich wykonywanych na nim bloków. Maksymalna liczba bloków na multiprocesor dla G8x/G9x wynosi osiem, a liczba wypaczeń to 24 (768 wątków na multiprocesor). W sumie najlepsze karty graficzne z serii Geforce 8 i 9 mogą jednocześnie przetwarzać do 12288 wątków. GeForce GTX 280 oparty na GT200 oferuje do 1024 wątków na multiprocesor, ma 10 klastrów trzech multiprocesorów przetwarzających do 30720 wątków. Znajomość tych ograniczeń pozwala na optymalizację algorytmów pod kątem dostępnych zasobów.

Pierwszym krokiem w przeniesieniu istniejącej aplikacji do CUDA jest jej profilowanie i identyfikowanie obszarów kodu, które są wąskimi gardłami spowalniającymi pracę. Jeśli wśród takich sekcji znajdują się odpowiednie do szybkiego wykonywania równoległego, funkcje te są przenoszone do rozszerzeń C i CUDA w celu wykonania na GPU. Program jest kompilowany przy użyciu kompilatora dostarczonego przez firmę Nvidia, który generuje kod zarówno dla CPU, jak i GPU. Kiedy program jest wykonywany, procesor wykonuje swoje części kodu, a GPU wykonuje kod CUDA z najcięższymi obliczeniami równoległymi. Ta część, zaprojektowana dla GPU, nazywa się jądrem (jądrem). Jądro definiuje operacje, które należy wykonać na danych.

Chip wideo odbiera rdzeń i tworzy kopie dla każdego elementu danych. Te kopie nazywane są wątkami. Strumień zawiera licznik, rejestry i stan. W przypadku dużych ilości danych, takich jak przetwarzanie obrazu, uruchamiane są miliony wątków. Wątki działają w grupach zwanych wypaczeniami po 32. Warpy są przypisane do działania na określonych multiprocesorach strumieniowych. Każdy multiprocesor składa się z ośmiu rdzeni - procesorów strumieniowych, które wykonują jedną instrukcję MAD na cykl zegara. Aby wykonać jedno 32-wątkowe warp, wymagane są cztery cykle wieloprocesorowe (mówimy o częstotliwości domeny shadera, która wynosi 1,5 GHz i więcej).

Wieloprocesor nie jest tradycyjnym procesorem wielordzeniowym, jest dobrze przystosowany do pracy wielowątkowej, obsługując jednocześnie do 32. W każdym cyklu zegara sprzęt wybiera, które z wypaczeń ma wykonać, i przełącza się z jednego na drugie bez przegrywanie cykli. Jeśli narysujemy analogię z procesorem centralnym, to tak, jakby wykonywać jednocześnie 32 programy i przełączać się między nimi w każdym cyklu zegara bez utraty przełącznika kontekstu. W rzeczywistości rdzenie procesora obsługują jednoczesne wykonywanie jednego programu i przełączanie się na inne z opóźnieniem setek cykli.

Model programowania CUDA

Ponownie, CUDA używa modelu obliczeń równoległych, w którym każdy z procesorów SIMD wykonuje równolegle tę samą instrukcję na różnych elementach danych. GPU jest urządzeniem obliczeniowym, koprocesorem (urządzeniem) dla procesora centralnego (hosta), który posiada własną pamięć i przetwarza równolegle dużą liczbę wątków. Jądro (kernel) to funkcja dla GPU, realizowana przez wątki (analogia do grafiki 3D - shader).

Powiedzieliśmy powyżej, że układ wideo różni się od procesora tym, że może przetwarzać dziesiątki tysięcy wątków jednocześnie, co zwykle dotyczy grafiki, która jest dobrze zrównoleglona. Każdy strumień jest skalarny, nie wymaga pakowania danych w 4-składnikowe wektory, co jest wygodniejsze w przypadku większości zadań. Liczba wątków logicznych i bloków wątków przewyższa liczbę fizycznych jednostek wykonawczych, co daje dobrą skalowalność dla całej gamy rozwiązań firmy.

Model programowania w CUDA zakłada grupowanie wątków. Wątki są łączone w bloki wątków - jednowymiarowe lub dwuwymiarowe siatki wątków oddziałujących ze sobą za pomocą pamięci współdzielonej i punktów synchronizacji. Program (jądro) jest wykonywany na siatce bloków wątków, patrz rysunek poniżej. Jednocześnie wykonywana jest jedna siatka. Każdy blok może mieć kształt jedno-, dwu- lub trójwymiarowy i może składać się z 512 wątków na obecnym sprzęcie.

Bloki wątków działają w małych grupach zwanych wypaczeniami, które mają rozmiar 32 wątków. Jest to minimalna ilość danych, którą można przetworzyć na wieloprocesorach. A ponieważ nie zawsze jest to wygodne, CUDA umożliwia pracę z blokami zawierającymi od 64 do 512 wątków.

Grupowanie bloków w siatki pozwala uciec od ograniczeń i zastosować jądro do większej liczby wątków w jednym wywołaniu. Pomaga również przy skalowaniu. Jeśli GPU nie ma wystarczających zasobów, wykona bloki sekwencyjnie. W przeciwnym razie bloki mogą być wykonywane równolegle, co jest ważne dla optymalnego rozłożenia pracy na chipach wideo na różnych poziomach, od mobilnych po zintegrowane.

Model pamięci CUDA

Model pamięci w CUDA wyróżnia się możliwością adresowania bajtowego, obsługą zarówno zbierania, jak i rozpraszania. Dla każdego procesora strumieniowego dostępna jest dość duża liczba rejestrów, do 1024 sztuk. Dostęp do nich jest bardzo szybki, można w nich przechowywać 32-bitowe liczby całkowite lub liczby zmiennoprzecinkowe.

Każdy wątek ma dostęp do następujących typów pamięci:

pamięć globalna- największa ilość pamięci dostępna dla wszystkich multiprocesorów na chipie wideo, wielkość waha się od 256 megabajtów do 1,5 gigabajta dla obecnych rozwiązań (i do 4 GB dla Tesli). Ma wysoką przepustowość, ponad 100 gigabajtów/s dla topowych rozwiązań Nvidii, ale bardzo duże opóźnienia rzędu kilkuset cykli. Nie można buforować, obsługuje ogólne instrukcje ładowania i przechowywania oraz zwykłe wskaźniki pamięci.

pamięć lokalna to niewielka ilość pamięci, do której dostęp ma tylko jeden procesor strumieniowy. Jest stosunkowo powolny - taki sam jak globalny.

Pamięć współdzielona to 16-kilobajtowy (w układach wideo obecnej architektury) blok pamięci ze współdzielonym dostępem dla wszystkich procesorów strumieniowych w multiprocesorze. Ta pamięć jest bardzo szybka, podobnie jak rejestry. Zapewnia interakcję z wątkiem, jest bezpośrednio zarządzany przez programistę i ma małe opóźnienia. Zalety pamięci współdzielonej: wykorzystanie w postaci pamięci podręcznej pierwszego poziomu zarządzanej przez programistę, zmniejszenie opóźnień w dostępie do danych przez jednostki wykonawcze (ALU), zmniejszenie liczby globalnych dostępów do pamięci.

Stała pamięć- 64-kilobajtowy obszar pamięci (taki sam dla obecnych procesorów graficznych), tylko do odczytu przez wszystkie procesory wieloprocesorowe. Jest buforowany przy 8 kilobajtach na wieloprocesor. Dość wolno - opóźnienie kilkuset cykli w przypadku braku niezbędnych danych w pamięci podręcznej.

pamięć tekstur- blok pamięci dostępny do odczytu przez wszystkie multiprocesory. Próbkowanie danych odbywa się za pomocą jednostek tekstury chipa wideo, dzięki czemu możliwość liniowej interpolacji danych jest zapewniona bez dodatkowych kosztów. 8 kilobajtów w pamięci podręcznej na wieloprocesor. Wolno jak globalnie - setki cykli opóźnienia w przypadku braku danych w pamięci podręcznej.

Naturalnie, pamięć globalna, lokalna, teksturowa i stała to fizycznie ta sama pamięć, znana jako lokalna pamięć wideo karty graficznej. Różnice dotyczą różnych algorytmów buforowania i modeli dostępu. Procesor może tylko aktualizować i wysyłać zapytania pamięć zewnętrzna: globalna, stała i tekstura.

Z tego, co zostało napisane powyżej, jasno wynika, że ​​CUDA implikuje specjalne podejście do programowania, nie całkiem takie samo, jak to przyjęte w programach dla CPU. Należy pamiętać o różnych typach pamięci, że pamięć lokalna i globalna nie są buforowane, a opóźnienia w dostępie do niej są znacznie większe niż w przypadku pamięci rejestrowanej, ponieważ fizycznie znajduje się ona w osobnych mikroukładach.

Typowy, ale nie obowiązkowy schemat rozwiązywania problemów:

• zadanie podzielone jest na podzadania;
• dane wejściowe są podzielone na bloki, które mieszczą się w pamięci współdzielonej;
• każdy blok jest przetwarzany przez blok wątku;
• podblok jest ładowany do pamięci współdzielonej z globalnej;
• odpowiednie obliczenia są wykonywane na danych w pamięci współdzielonej;
• wyniki są kopiowane z pamięci współdzielonej z powrotem do globalnej.

Środowisko programistyczne

CUDA zawiera biblioteki wykonawcze:

• wspólna część, która zapewnia wbudowane typy wektorów i podzbiory wywołań RTL obsługiwanych przez CPU i GPU;
• składnik procesora do sterowania jednym lub większą liczbą procesorów graficznych;
• Składnik GPU, który zapewnia funkcje specyficzne dla GPU.

Główny proces aplikacji CUDA działa na ogólnym procesorze (hoście), uruchamia wiele kopii procesów jądra na karcie graficznej. Kod procesora wykonuje następujące czynności: inicjuje GPU, przydziela pamięć na karcie graficznej i systemie, kopiuje stałe do pamięci karty graficznej, uruchamia kilka kopii procesów jądra na karcie graficznej, kopiuje wynik z pamięć wideo, zwalnia pamięć i wychodzi.

Jako przykład do zrozumienia, oto kod procesora do dodawania wektorów przedstawiony w CUDA:

Funkcje wykonywane przez układ wideo mają następujące ograniczenia: brak rekurencji, brak zmiennych statycznych wewnątrz funkcji oraz brak zmiennej liczby argumentów. Obsługiwane są dwa typy zarządzania pamięcią: pamięć liniowa dostępna przez 32-bitowe wskaźniki oraz tablice CUDA dostępne tylko przez funkcje pobierania tekstur.

Programy CUDA mogą wchodzić w interakcje z graficznymi interfejsami API: renderować dane generowane w programie, odczytywać wyniki renderowania i przetwarzać je za pomocą narzędzi CUDA (na przykład podczas implementacji filtrów przetwarzania końcowego). W tym celu zasoby graficznego interfejsu API mogą być mapowane (uzyskując adres zasobu) do globalnej przestrzeni pamięci CUDA. Obsługiwane są następujące typy graficznych zasobów API: Obiekty buforowe (PBO / VBO) w OpenGL, bufory wierzchołków i tekstury (2D, 3D i cubemaps) Direct3D9.

Kroki kompilacji aplikacji CUDA:

Pliki z kodem źródłowym CUDA C są kompilowane przy użyciu programu NVCC, który otacza inne narzędzia i wywołuje je: cudacc, g++, cl, itp. NVCC generuje: kod procesora, który jest kompilowany wraz z resztą aplikacji napisaną w czystym C oraz Kod obiektowy PTX dla układu wideo. Pliki wykonywalne z kodem CUDA koniecznie wymagają obecności biblioteki wykonawczej CUDA (cudart) i podstawowej biblioteki CUDA (cuda).

Optymalizacja programów na CUDA

Oczywiście w ramach artykułu przeglądowego nie można rozważać poważnych problemów optymalizacyjnych w programowaniu CUDA. Dlatego pokrótce porozmawiamy o podstawowych rzeczach. Aby efektywnie wykorzystać możliwości CUDA, trzeba zapomnieć o zwykłych metodach pisania programów dla CPU i korzystać z tych algorytmów, które są dobrze zrównoleglone dla tysięcy wątków. Ważne jest również znalezienie optymalnego miejsca do przechowywania danych (rejestry, pamięć współdzielona itp.), zminimalizowanie transferu danych między procesorem a GPU oraz użycie buforowania.

Ogólnie rzecz biorąc, podczas optymalizacji programu CUDA należy starać się osiągnąć optymalną równowagę między rozmiarem a liczbą bloków. Więcej wątków w bloku zmniejszy wpływ opóźnienia pamięci, ale zmniejszy również dostępną liczbę rejestrów. Ponadto blok 512 wątków jest nieefektywny, sama Nvidia zaleca używanie bloków 128 lub 256 wątków jako wartości kompromisowej w celu uzyskania optymalnego opóźnienia i liczby rejestrów.

Wśród głównych punktów optymalizacji programów CUDA: jak najbardziej aktywne wykorzystanie pamięci współdzielonej, ponieważ jest ona znacznie szybsza niż globalna pamięć wideo karty graficznej; odczyty i zapisy z pamięci globalnej powinny być łączone, gdy tylko jest to możliwe. Aby to zrobić, musisz użyć specjalnych typów danych do jednoczesnego odczytu i zapisu 32/64/128 bitów danych w jednej operacji. Jeśli operacje odczytu są trudne do scalenia, możesz spróbować użyć pobierania tekstur.

wnioski

Architektura sprzętowa i programowa prezentowana przez Nvidię do obliczeń na chipach wideo CUDA jest dobrze dostosowana do rozwiązywania szerokiego zakresu zadań o dużej równoległości. CUDA działa na dużej liczbie układów wideo Nvidii i poprawia model programowania GPU, znacznie go upraszczając i dodając wiele funkcji, takich jak pamięć współdzielona, ​​możliwość synchronizacji wątków, obliczenia o podwójnej precyzji i operacje na liczbach całkowitych.

CUDA to technologia dostępna dla każdego programisty, może być używana przez każdego programistę znającego język C. Musisz tylko przyzwyczaić się do innego paradygmatu programowania związanego z przetwarzaniem równoległym. Ale jeśli algorytm jest w zasadzie dobrze zrównoleglony, to nauka i czas spędzony na programowaniu CUDA zwrócą się w wielu rozmiarach.

Jest prawdopodobne, że ze względu na powszechne stosowanie kart graficznych na świecie rozwój obliczeń równoległych na GPU znacznie wpłynie na branżę obliczeń o wysokiej wydajności. Możliwości te wzbudziły już duże zainteresowanie środowisk naukowych i nie tylko. W końcu potencjał przyspieszania algorytmów, które dobrze nadają się do zrównoleglania (na niedrogim sprzęcie, co nie mniej ważne) od razu, dziesiątki razy nie jest tak powszechny.

Procesory ogólnego przeznaczenia rozwijają się dość wolno, nie mają takich skoków wydajności. W rzeczywistości, chociaż brzmi to zbyt głośno, każdy, kto potrzebuje szybkich komputerów, może teraz kupić niedrogi osobisty superkomputer na swoim biurku, czasami nawet bez inwestycji dodatkowe środki, ponieważ karty graficzne Nvidii są szeroko stosowane. Nie wspominając o wzroście wydajności pod względem GFLOPS/$ i GFLOPS/W, które tak bardzo kochają producenci GPU.

Przyszłość wielu obliczeń jest wyraźnie w równoległych algorytmach, prawie wszystkie nowe rozwiązania i inicjatywy są skierowane w tym kierunku. Póki co jednak trwa opracowywanie nowych paradygmatów etap początkowy, musisz ręcznie tworzyć wątki i planować dostęp do pamięci, co sprawia, że ​​rzeczy są trudniejsze niż normalnie podczas programowania. Ale technologia CUDA zrobiła krok we właściwym kierunku i wyraźnie wygląda na udane rozwiązanie, zwłaszcza jeśli Nvidia zdoła przekonać deweloperów jak najwięcej o jej zaletach i perspektywach.

Ale oczywiście GPU nie zastąpią procesorów. W obecnej formie nie są do tego przeznaczone. Teraz, gdy chipy wideo stopniowo zbliżają się do procesora, stając się coraz bardziej uniwersalne (obliczenia zmiennoprzecinkowe pojedynczej i podwójnej precyzji, obliczenia na liczbach całkowitych), więc procesory stają się coraz bardziej „równoległe”, zdobywając dużą liczbę rdzeni, wielowątkowość technologii, nie wspominając o pojawieniu się bloków SIMD i projektów heterogenicznych procesorów. Najprawdopodobniej w przyszłości GPU i CPU po prostu się połączą. Wiadomo, że nad podobnymi projektami pracuje wiele firm, w tym Intel i AMD. I nie ma znaczenia, czy procesor graficzny jest zużywany przez procesor, czy odwrotnie.

W artykule rozmawialiśmy głównie o zaletach CUDA. Ale w maści jest też mucha. Jedną z niewielu wad CUDA jest jej słaba przenośność. Ta architektura działa tylko na chipach wideo tej firmy, a nie na wszystkich, ale zaczynając od serii Geforce 8 i 9 oraz odpowiednich Quadro i Tesli. Tak, na świecie jest wiele takich rozwiązań, Nvidia podaje liczbę 90 milionów chipów wideo zgodnych z CUDA. To jest po prostu świetne, ale konkurenci oferują własne rozwiązania, które różnią się od CUDA. Tak więc AMD ma Stream Computing, Intel będzie miał Ct w przyszłości.

Która z technologii zwycięży, stanie się powszechna i będzie żyć dłużej niż pozostałe - tylko czas pokaże. Ale CUDA ma duże szanse, ponieważ w porównaniu na przykład do Stream Computing zapewnia bardziej rozwinięte i łatwe w użyciu środowisko programistyczne w zwykłym języku C. Być może osoba trzecia pomoże w ustaleniu wydając jakieś wspólna decyzja. Na przykład w następnej aktualizacji DirectX w wersji 11 Microsoft obiecał shadery obliczeniowe, które mogą stać się rodzajem przeciętnego rozwiązania, które odpowiada każdemu lub prawie każdemu.

Sądząc po wstępnych danych, ten nowy rodzaj shadera wiele zapożycza z modelu CUDA. A dzięki programowaniu w tym środowisku już teraz możesz uzyskać natychmiastowe korzyści i niezbędne umiejętności na przyszłość. Z punktu widzenia obliczeń o wysokiej wydajności, DirectX ma również wyraźną wadę, jaką jest słaba przenośność, ponieważ API jest ograniczone do platformy Windows. Jednak opracowywany jest inny standard - otwarta wieloplatformowa inicjatywa OpenCL, którą wspiera większość firm, w tym Nvidia, AMD, Intel, IBM i wiele innych.

Nie zapominaj, że w następnym artykule CUDA poznasz konkretne praktyczne zastosowania obliczeń naukowych i innych niegraficznych, wykonywanych przez programistów na całym świecie za pomocą Nvidia CUDA.

Wróćmy do historii – cofnijmy się do 2003 roku, kiedy Intel i AMD ścigały się wspólnie o najmocniejszy procesor. W ciągu zaledwie kilku lat, w wyniku tego wyścigu, taktowanie zegarów znacznie wzrosło, zwłaszcza po wydaniu Intela Pentium 4.

Ale wyścig szybko zbliżał się do limitu. Po fali ogromnych wzrostów częstotliwości taktowania (w latach 2001-2003 taktowanie Pentium 4 podwoiło się z 1,5 do 3 GHz), użytkownicy musieli zadowolić się dziesiątkami gigaherca, które producenci byli w stanie wycisnąć (od 2003 do 2005). , taktowanie zwiększyło się z zaledwie 3 do 3,8 GHz).

Nawet architektury zoptymalizowane pod kątem wysokich częstotliwości taktowania, jak Prescott, zaczęły doświadczać trudności, tym razem nie tylko produkcyjnych. Twórcy chipów właśnie natknęli się na prawa fizyki. Niektórzy analitycy przewidywali nawet, że prawo Moore'a przestanie działać. Ale tak się nie stało. Pierwotne znaczenie prawa jest często błędnie przedstawiane, ale odnosi się ono do liczby tranzystorów na powierzchni rdzenia krzemowego. Przez długi czas wzrostowi liczby tranzystorów w CPU towarzyszył odpowiedni wzrost wydajności - co prowadziło do zniekształcenia znaczenia. Ale potem sytuacja się skomplikowała. Projektanci architektury procesora podeszli do prawa redukcji wzmocnienia: liczba tranzystorów, które trzeba było dodać, aby uzyskać pożądany wzrost wydajności, stawała się coraz większa, prowadząc do ślepego zaułka.

Podczas gdy producenci procesorów wyrywają sobie włosy z włosów, próbując znaleźć rozwiązanie swoich problemów, producenci procesorów graficznych nadal odnoszą znaczne korzyści z dobrodziejstw prawa Moore'a.

Dlaczego nie znaleźli się w tym samym ślepym zaułku, co projektanci architektury CPU? Powód jest bardzo prosty: procesory są zaprojektowane tak, aby uzyskać najlepszą wydajność w strumieniu instrukcji, które przetwarzają różne dane (zarówno liczby całkowite, jak i liczby zmiennoprzecinkowe), wykonują losowy dostęp do pamięci i tak dalej. Do tej pory programiści starali się zapewnić większą równoległość instrukcji - to znaczy wykonywać równolegle jak najwięcej instrukcji. Tak więc na przykład w Pentium pojawiło się wykonanie superskalarne, gdy w określonych warunkach możliwe było wykonanie dwóch instrukcji na zegar. Pentium Pro otrzymał wykonanie instrukcji poza kolejnością, co pozwoliło zoptymalizować wydajność jednostek obliczeniowych. Problem polega na tym, że równoległe wykonywanie sekwencyjnego strumienia instrukcji ma oczywiste ograniczenia, więc ślepe zwiększanie liczby jednostek obliczeniowych nie daje zysku, ponieważ przez większość czasu będą one nadal bezczynne.

Wręcz przeciwnie, praca GPU jest stosunkowo prosta. Polega na wzięciu grupy wielokątów z jednej strony i wygenerowaniu grupy pikseli z drugiej. Wielokąty i piksele są od siebie niezależne, dzięki czemu można je przetwarzać równolegle. Tym samym w GPU można przeznaczyć dużą część kryształu na jednostki obliczeniowe, które w przeciwieństwie do CPU będą faktycznie wykorzystywane.

GPU różni się od CPU nie tylko tym. Dostęp do pamięci w GPU jest mocno sprzężony - jeśli odczytany zostanie teksel, to po kilku cyklach zostanie odczytany teksel sąsiedni; gdy piksel zostanie zapisany, sąsiedni zostanie zapisany po kilku cyklach. Inteligentnie organizując pamięć, można uzyskać wydajność zbliżoną do teoretycznej przepustowości. Oznacza to, że GPU, w przeciwieństwie do CPU, nie wymaga ogromnej pamięci podręcznej, ponieważ jego rolą jest przyspieszenie operacji teksturowania. Wystarczy kilka kilobajtów zawierających kilka tekseli używanych w filtrach dwu- i trójliniowych.

Niech żyje GeForce FX!

Dwa światy przez długi czas pozostawały rozdzielone. Używaliśmy procesora (lub nawet wielu procesorów) do zadań biurowych i aplikacji internetowych, a GPU doskonale nadawał się tylko do przyspieszenia renderowania. Ale jedna funkcja zmieniła wszystko: a mianowicie pojawienie się programowalnych procesorów graficznych. Początkowo procesory nie miały się czego obawiać. Pierwsze tak zwane programowalne GPU (NV20 i R200) ​​nie stanowiły większego zagrożenia. Liczba instrukcji w programie pozostała ograniczona do około 10, pracowały na bardzo egzotycznych typach danych, takich jak 9- lub 12-bitowe liczby stałoprzecinkowe.

Ale prawo Moore'a znów pokazało się z najlepszej strony. Wzrost liczby tranzystorów nie tylko zwiększył liczbę jednostek obliczeniowych, ale także poprawił ich elastyczność. Pojawienie się NV30 można uznać za znaczący krok naprzód z kilku powodów. Oczywiście gracze nie przepadali za kartami NV30, ale nowe procesory graficzne zaczęły opierać się na dwóch funkcjach, które miały zmienić postrzeganie procesorów graficznych jako czegoś więcej niż tylko akceleratorów graficznych.

• Obsługa obliczeń zmiennoprzecinkowych o pojedynczej precyzji (nawet jeśli nie jest zgodny ze standardem IEEE754);
• obsługa ponad tysiąca instrukcji.

Stworzyliśmy więc wszystkie warunki, które mogą przyciągnąć pionierskich naukowców, którzy zawsze szukają dodatkowej mocy obliczeniowej.

Pomysł wykorzystania akceleratorów graficznych do obliczeń matematycznych nie jest nowy. Pierwsze próby podjęto w latach 90. ubiegłego wieku. Oczywiście były one bardzo prymitywne – w większości ograniczały się do korzystania z niektórych funkcji sprzętowych, takich jak rasteryzacja i bufory Z w celu przyspieszenia zadań, takich jak wyszukiwanie trasy lub wyjście Diagramy Woronoja .

W 2003 roku, wraz z pojawieniem się ewoluujących shaderów, osiągnięto nowy pasek - tym razem wykonujący obliczenia macierzowe. W tym roku cała sekcja SIGGRAPH („Obliczenia na GPU”) została poświęcona nowemu obszarowi IT. Ta wczesna inicjatywa została nazwana GPGPU (General-Purpose Computation na GPU). I pojawienie się .

Aby zrozumieć rolę BrookGPU, musisz zrozumieć, jak wszystko się wydarzyło przed jego pojawieniem się. Jedynym sposobem na uzyskanie zasobów GPU w 2003 roku było użycie jednego z dwóch graficznych API - Direct3D lub OpenGL. W związku z tym programiści, którzy chcieli uzyskać moc GPU do swoich obliczeń, musieli polegać na dwóch wspomnianych interfejsach API. Problem w tym, że nie zawsze byli ekspertami w programowaniu kart graficznych, a to bardzo utrudniało dostęp do technologii. Jeśli programiści 3D operują na shaderach, teksturach i fragmentach, to specjaliści z zakresu programowania równoległego polegają na wątkach, rdzeniach, rozpraszaniach itp. Dlatego na początku konieczne było wyciągnięcie analogii między dwoma światami.

• strumień to strumień elementów tego samego typu, w GPU może być reprezentowany przez teksturę. W zasadzie w programowaniu klasycznym istnieje taki analog, jak tablica.
• Jądro- funkcja, która zostanie zastosowana niezależnie do każdego elementu strumienia; jest odpowiednikiem Pixel Shadera. W programowaniu klasycznym można podać analogię do cyklu - dotyczy on dużej liczby elementów.
• Aby odczytać wyniki zastosowania jądra do strumienia, należy utworzyć teksturę. Nie ma odpowiednika na procesorze, ponieważ jest pełny dostęp do pamięci.
• Lokalizacja w pamięci, do której ma zostać zapisany (w operacjach rozpraszania/rozproszenia), jest kontrolowana przez Vertex Shader, ponieważ Pixel Shader nie może zmienić współrzędnych przetwarzanego piksela.

Jak widać, nawet biorąc pod uwagę powyższe analogie, zadanie nie wygląda na proste. I Brook przyszedł na ratunek. Ta nazwa odnosi się do rozszerzeń języka C („C ze strumieniami”, „C ze strumieniami”), jak nazywali je programiści ze Stanford. W istocie zadaniem Brooka było ukrycie wszystkich komponentów API 3D przed programistą, co umożliwiło przedstawienie GPU jako koprocesora do obliczeń równoległych. W tym celu kompilator Brook przetworzył plik .br z kodem C++ i rozszerzeniami, a następnie wygenerował kod C++, który został połączony z biblioteką obsługującą różne wyjścia (DirectX, OpenGL ARB, OpenGL NV3x, x86).

Brook ma kilka zasług, z których pierwszym jest wyciągnięcie GPGPU z cienia, aby technologia była widoczna dla ogółu społeczeństwa. Choć po ogłoszeniu projektu wiele serwisów IT było zbyt optymistycznych, że premiera Brooka poddaje w wątpliwość istnienie procesorów, które niedługo zostaną zastąpione mocniejszymi GPU. Ale, jak widzimy, nawet po pięciu latach tak się nie stało. Szczerze mówiąc, sądzimy, że to się nigdy nie zdarzy. Z drugiej strony patrząc na udaną ewolucję procesorów, które są coraz bardziej zorientowane na równoległość (więcej rdzeni, technologia wielowątkowości SMT, rozbudowa bloków SIMD), a także GPU, które wręcz przeciwnie, stają się coraz bardziej uniwersalne (obsługa do obliczeń zmiennoprzecinkowych) pojedyncza precyzja, obliczenia na liczbach całkowitych, obsługa obliczeń podwójnej precyzji), wygląda na to, że GPU i CPU wkrótce po prostu się połączą. Co się wtedy stanie? Czy procesory graficzne zostaną wchłonięte przez procesory, jak miało to miejsce w przypadku koprocesorów matematycznych? Całkiem możliwe. Intel i AMD pracują dziś nad podobnymi projektami. Ale wiele może się jeszcze zmienić.

Wróćmy jednak do naszego tematu. Atutem Brooka było spopularyzowanie koncepcji GPGPU, znacznie uprościło to dostęp do zasobów GPU, co pozwoliło coraz większej liczbie użytkowników opanować nowy model programowania. Z drugiej strony, pomimo wszystkich zalet Brooka, wciąż była długa droga do wykorzystania zasobów GPU do obliczeń.

Jeden z problemów jest związany z różnymi poziomami abstrakcji, a także w szczególności z nadmiernym dodatkowym obciążeniem tworzonym przez 3D API, co może być dość zauważalne. Ale poważniejszy można uznać za problem ze zgodnością, z którym twórcy Brooka nie mogli nic zrobić. Między producentami procesorów graficznych panuje ostra konkurencja, dlatego często optymalizują oni swoje sterowniki. Jeśli takie optymalizacje są w większości dobre dla graczy, mogą w jednej chwili wyeliminować kompatybilność z Brook. Dlatego trudno sobie wyobrazić wykorzystanie tego API w kodzie przemysłowym, który gdzieś zadziała. I przez długi czas Brook pozostał grupą badaczy-amatorów i programistów.

Jednak sukces Brooka wystarczył, aby przyciągnąć uwagę ATI i Nvidii, które były zainteresowane taką inicjatywą, ponieważ mogłaby poszerzyć rynek, otwierając nowy ważny sektor dla firm.

Naukowcy początkowo zaangażowani w projekt Brook szybko dołączyli do zespołów programistycznych w Santa Clara, aby przedstawić globalną strategię rozwoju nowego rynku. Pomysł polegał na stworzeniu połączenia sprzętu i oprogramowania odpowiedniego do zadań GPGPU. Ponieważ twórcy nVidii znają wszystkie sekrety swoich procesorów graficznych, można było nie polegać na graficznym API, ale komunikować się z procesorem graficznym za pośrednictwem sterownika. Chociaż oczywiście wiąże się to z własnymi problemami. Tak więc zespół programistów CUDA (Compute Unified Device Architecture) stworzył zestaw warstw oprogramowania do pracy z GPU.

Jak widać na diagramie, CUDA udostępnia dwa API.

• API wysokiego poziomu: CUDA Runtime API;
• API niskiego poziomu: API sterownika CUDA.

Ponieważ interfejs API wysokiego poziomu jest zaimplementowany nad interfejsem niskiego poziomu, każde wywołanie funkcji środowiska wykonawczego jest dzielone na prostsze instrukcje, które są przetwarzane przez interfejs API sterownika. Zauważ, że te dwa API wykluczają się wzajemnie: programista może używać jednego lub drugiego API, ale nie jest możliwe łączenie wywołań funkcji z dwóch API. Ogólnie termin „API wysokiego poziomu” jest względny. Nawet Runtime API jest taki, że wielu uzna go za niskopoziomowy; jednak nadal zapewnia funkcje, które są bardzo wygodne do inicjowania lub zarządzania kontekstem. Ale nie oczekuj szczególnie wysokiego poziomu abstrakcji - nadal musisz mieć dobrą wiedzę na temat procesorów graficznych nVidia i ich działania.

Driver API jest jeszcze trudniejszy w obsłudze; potrzebujesz więcej wysiłku, aby uruchomić przetwarzanie GPU. Z drugiej strony, niskopoziomowe API jest bardziej elastyczne, dając programiście większą kontrolę w razie potrzeby. Dwa interfejsy API mogą pracować z zasobami OpenGL lub Direct3D (na dzień dzisiejszy tylko dziewiąta wersja). Zaleta tej funkcji jest oczywista - CUDA może być wykorzystana do tworzenia zasobów (geometrii, tekstur proceduralnych itp.), które można przekazać do API graficznego lub odwrotnie, API 3D może być wykorzystane do wysyłania wyników renderowania do CUDA program, który z kolei wykona przetwarzanie końcowe. Przykładów takich interakcji jest wiele, a zaletą jest to, że zasoby są nadal przechowywane w pamięci GPU, nie muszą być przesyłane przez magistralę PCI Express, która wciąż jest wąskim gardłem.

Należy jednak zauważyć, że współdzielenie zasobów w pamięci wideo nie zawsze jest idealne i może prowadzić do pewnych „bólów głowy”. Na przykład przy zmianie rozdzielczości lub głębi kolorów dane graficzne mają pierwszeństwo. Dlatego jeśli konieczne jest zwiększenie zasobów w buforze ramki, to sterownik z łatwością zrobi to kosztem zasobów aplikacji CUDA, które po prostu zawieszą się z błędem. Oczywiście niezbyt elegancka, ale taka sytuacja nie powinna zdarzać się zbyt często. A odkąd zaczęliśmy mówić o wadach: jeśli chcesz używać wielu procesorów graficznych do aplikacji CUDA, musisz najpierw wyłączyć tryb SLI, w przeciwnym razie aplikacje CUDA będą mogły „widzieć” tylko jeden procesor graficzny.

Wreszcie trzeci poziom oprogramowania jest przyznawany bibliotekom - a dokładnie dwa.

• CUBLAS, który zawiera niezbędne bloki do obliczania algebry liniowej na GPU;
• CUFFT, który wspiera obliczanie transformat Fouriera, to algorytm szeroko stosowany w dziedzinie przetwarzania sygnałów.

Zanim zagłębimy się w CUDA, zdefiniujmy kilka terminów rozproszonych w dokumentacji nVidii. Firma wybrała bardzo specyficzną terminologię, do której trudno się przyzwyczaić. Przede wszystkim zauważamy, że wątek w CUDA nie ma nawet takiego samego znaczenia jak wątek procesora, ani nie jest odpowiednikiem wątku w naszych artykułach dotyczących GPU. Wątek GPU w tym przypadku to zestaw podstawowy dane do przetwarzania. W przeciwieństwie do wątków procesora, wątki CUDA są bardzo „lekkie”, co oznacza, że ​​przełączanie kontekstu między dwoma wątkami w żadnym wypadku nie jest operacją intensywnie wykorzystującą zasoby.

Drugi termin często spotykany w dokumentacji CUDA to: osnowa. Tutaj nie ma zamieszania, ponieważ nie ma odpowiednika w języku rosyjskim (chyba że jesteś fanem gier Start Trek lub Warhammer). W rzeczywistości termin ten został zaczerpnięty z przemysłu tekstylnego, gdzie przędza na wątek jest przeciągana przez przędzę osnowową naciągniętą na krośnie. Warp w CUDA to grupa 32 wątków i jest to minimalna ilość danych przetwarzanych w sposób SIMD w wieloprocesorach CUDA.

Ale taka „ziarnistość” nie zawsze jest wygodna dla programisty. Dlatego w CUDA, zamiast bezpośrednio pracować z wypaczeniami, możesz pracować z blok, zawierający od 64 do 512 wątków.

Wreszcie te bloki są ze sobą połączone w siatki. Zaletą tego grupowania jest to, że liczba bloków przetwarzanych przez GPU w tym samym czasie jest ściśle powiązana z zasobami sprzętowymi, co zobaczymy poniżej. Grupowanie bloków w siatki pozwala całkowicie oderwać się od tego ograniczenia i zastosować jądro/jądro do większej liczby wątków w jednym wywołaniu, bez myślenia o stałych zasobach. Za to wszystko odpowiadają biblioteki CUDA. Ponadto model ten dobrze się skaluje. Jeśli GPU ma mało zasobów, bloki będą wykonywane sekwencyjnie. Jeśli liczba procesorów obliczeniowych jest duża, bloki mogą być wykonywane równolegle. Oznacza to, że ten sam kod może działać zarówno na procesorach graficznych klasy podstawowej, jak i na topowych, a nawet przyszłych modelach.

W interfejsie API CUDA jest jeszcze kilka terminów oznaczających procesor ( gospodarz/gospodarz) i GPU ( urządzenie/urządzenie). Jeśli ten mały wstęp Cię nie przestraszył, to czas lepiej poznać CUDA.

Jeśli regularnie czytasz Tom's Hardware Guide, architektura najnowszych procesorów graficznych nVidia jest Ci znana. Jeśli nie, zalecamy przeczytanie artykułu „ nVidia GeForce GTX 260 i 280: nowa generacja kart graficznych Jeśli chodzi o CUDA, nVidia prezentuje architekturę nieco inaczej, pokazując niektóre szczegóły, które wcześniej były ukryte.

Jak widać na powyższej ilustracji, rdzeń shadera nVidia składa się z kilku klastrów procesorów tekstur. (Klaster procesorów tekstur, TPC). Na przykład 8800 GTX używał ośmiu klastrów, 8800 GTS używał sześciu i tak dalej. Każdy klaster zasadniczo składa się z jednostki tekstury i dwóch multiprocesory strumieniowe. Te ostatnie obejmują początek potoku (front end), który odczytuje i dekoduje instrukcje, a także wysyła je do wykonania, oraz koniec potoku (back end), składający się z ośmiu urządzeń obliczeniowych i dwóch urządzeń superfunkcyjnych. SFU (jednostka superfunkcyjna), gdzie instrukcje są wykonywane zgodnie z zasadą SIMD, to znaczy jedna instrukcja jest stosowana do wszystkich wątków w osnowie. nVidia nazywa ten sposób robienia tego SIMT(pojedyncza instrukcja wiele wątków, jedna instrukcja, wiele wątków). Należy zauważyć, że koniec rurociągu działa z dwukrotnie większą częstotliwością niż jego początek. W praktyce oznacza to, że część wygląda na dwa razy „szerszą” niż jest w rzeczywistości (tj. jak 16-kanałowy blok SIMD zamiast ośmiokanałowego). Multiprocesory strumieniujące działają w następujący sposób: w każdym cyklu, początek potoku wybiera odkształcenie gotowe do wykonania i rozpoczyna wykonywanie instrukcji. Aby instrukcja miała zastosowanie do wszystkich 32 wątków w osnowie, koniec potoku zająłby cztery cykle zegara, ale ponieważ przebiega on z dwukrotnie większą częstotliwością niż początek, zajęłoby to tylko dwa cykle zegara (pod względem początku rurociąg). Dlatego, aby początek potoku nie pozostawał bezczynny przez cykl, a sprzęt był maksymalnie obciążony, w idealnym przypadku można naprzemiennie instrukcje w każdym cyklu - klasyczną instrukcję w jednym cyklu i instrukcję dla SFU - w innym .

Każdy multiprocesor ma określony zestaw zasobów, które warto zrozumieć. Istnieje niewielki obszar pamięci zwany „Współdzielona pamięć”, 16 KB na wieloprocesor. Nie jest to bynajmniej pamięć podręczna: programista może z niej korzystać według własnego uznania. Oznacza to, że mamy coś zbliżonego do lokalnego sklepu SPU na procesorach Cell. Ten szczegół jest dość interesujący, ponieważ podkreśla, że ​​CUDA to połączenie technologii oprogramowania i sprzętu. Ten obszar pamięci nie jest wykorzystywany do shaderów pikseli, jak Nvidia sprytnie wskazuje „nie lubimy, gdy piksele rozmawiają ze sobą”.

Ten obszar pamięci otwiera możliwość wymiany informacji między wątkami. w jednym bloku. Należy podkreślić to ograniczenie: gwarantuje się, że wszystkie wątki w bloku zostaną wykonane przez jeden wieloprocesor. Wręcz przeciwnie, wiązanie bloków z różnymi multiprocesorami nie jest w ogóle przewidziane, a dwa wątki z różnych bloków nie mogą wymieniać między sobą informacji podczas wykonywania. Oznacza to, że korzystanie z pamięci współdzielonej nie jest takie proste. Jednak pamięć współdzielona jest nadal uzasadniona, z wyjątkiem sytuacji, gdy wiele wątków próbuje uzyskać dostęp do tego samego banku pamięci, powodując konflikt. W innych sytuacjach dostęp do pamięci współdzielonej jest tak szybki, jak dostęp do rejestru.

Pamięć współdzielona nie jest jedyną, do której mają dostęp wieloprocesory. Mogą korzystać z pamięci wideo, ale z mniejszą przepustowością i większym opóźnieniem. Dlatego, aby zmniejszyć częstotliwość dostępu do tej pamięci, nVidia wyposażyła multiprocesory w pamięć podręczną (około 8 KB na multiprocesor), która przechowuje stałe i tekstury.

Multiprocesor ma 8192 rejestrów, które są wspólne dla wszystkich wątków wszystkich bloków aktywnych na multiprocesorze. Liczba aktywnych bloków na multiprocesor nie może przekraczać ośmiu, a liczba aktywnych wypaczeń jest ograniczona do 24 (768 wątków). Dlatego 8800 GTX może przetwarzać do 12 288 wątków jednocześnie. Warto wspomnieć o wszystkich tych ograniczeniach, ponieważ pozwalają one na optymalizację algorytmu w oparciu o dostępne zasoby.

Optymalizacja programu CUDA polega zatem na uzyskaniu optymalnej równowagi między liczbą bloków a ich rozmiarem. Więcej wątków na blok byłoby pomocne w zmniejszeniu opóźnienia pamięci, ale liczba rejestrów dostępnych na wątek również byłaby zmniejszona. Co więcej, blok 512 wątków byłby nieefektywny, ponieważ tylko jeden blok mógłby być aktywny na wieloprocesorze, powodując utratę 256 wątków. Dlatego nVidia zaleca używanie bloków 128 lub 256 wątków, co daje najlepszy kompromis między mniejszą latencją a liczbą rejestrów dla większości rdzeni/jąder.

Z programowego punktu widzenia CUDA składa się z zestawu rozszerzeń do języka C, który przypomina BrookGPU, a także z kilku specyficznych wywołań API. Wśród rozszerzeń znajdują się specyfikatory typu związane z funkcjami i zmiennymi. Ważne jest, aby zapamiętać słowo kluczowe __światowy__, który podany przed funkcją wskazuje, że ta ostatnia odnosi się do jądra / jądra - ta funkcja zostanie wywołana przez CPU i zostanie wykonana na GPU. Prefiks __urządzenie__ określa, że ​​funkcja będzie wykonywana na GPU (którą CUDA przy okazji nazywa „urządzeniem/urządzeniem”), ale można ją wywołać tylko z GPU (innymi słowy, z innej funkcji __device__ lub z funkcji __global__). Wreszcie przedrostek __gospodarz__ opcjonalny, oznacza funkcję, która jest wywoływana przez CPU i wykonywana przez CPU - innymi słowy zwykła funkcja.

Istnieje wiele ograniczeń związanych z funkcjami __device__ i __global__: nie mogą one być rekurencyjne (czyli wywoływać siebie) i nie mogą mieć zmiennej liczby argumentów. Wreszcie, ponieważ funkcje __device__ znajdują się w przestrzeni pamięci GPU, ma sens, że nie można pobrać ich adresu. Zmienne mają również szereg kwalifikatorów, które wskazują lokalizację pamięci, w której będą przechowywane. Zmienna z prefiksem __wspólny__ oznacza, że ​​będzie przechowywany w pamięci współdzielonej multiprocesora strumieniowego. Wywołanie funkcji __global__ jest nieco inne. Chodzi o to, że podczas wywoływania musisz ustawić konfigurację wykonania - a dokładniej rozmiar siatki / siatki, do której zostanie zastosowane jądro / jądro, a także rozmiar każdego bloku. Weźmy na przykład jądro z następującą sygnaturą.

Będzie się nazywać jako

gdzie Dg to rozmiar siatki, a Db to rozmiar bloku. Te dwie zmienne odnoszą się do nowego typu wektora wprowadzonego w CUDA.

API CUDA zawiera funkcje do pracy z pamięcią w VRAM: cudaMalloc do przydzielania pamięci, cudaFree do zwalniania i cudaMemcpy do kopiowania pamięci między RAM i VRAM i na odwrót.

Tę recenzję zakończymy bardzo ciekawym sposobem kompilacji programu CUDA: kompilacja odbywa się w kilku krokach. Po pierwsze, kod specyficzny dla procesora jest wyodrębniany i przekazywany do standardowego kompilatora. Kod przeznaczony dla GPU jest najpierw konwertowany na język pośredni PTX. Jest podobny do języka asemblera i pozwala na badanie kodu w poszukiwaniu potencjalnie nieefektywnych sekcji. Ostatnią fazą jest przetłumaczenie języka pośredniego na instrukcje specyficzne dla GPU i utworzenie pliku binarnego.

Przeglądanie dokumentacji nVidii sprawia, że ​​chcę w tym tygodniu wypróbować CUDA. Rzeczywiście, co może być lepszego niż ocena API poprzez stworzenie własnego programu? Wtedy większość problemów powinna wyjść na jaw, nawet jeśli na papierze wszystko wygląda idealnie. Ponadto praktyka najlepiej pokaże, jak dobrze rozumiesz wszystkie zasady przedstawione w dokumentacji CUDA.

Zanurz się w takim projekcie dość łatwo. Obecnie do pobrania dostępna jest duża liczba bezpłatnych, ale wysokiej jakości narzędzi. Do naszego testu użyliśmy Visual C++ Express 2005, który ma wszystko, czego potrzebujesz. Najtrudniejsze było znalezienie programu, którego przeniesienie na GPU nie zajęło wielu tygodni, ale był na tyle interesujący, że nasze wysiłki nie poszły na marne. Na koniec wybraliśmy fragment kodu, który pobiera mapę wysokości i oblicza odpowiednią mapę normalnych. Nie będziemy zagłębiać się w tę funkcję szczegółowo, ponieważ nie jest ona interesująca w tym artykule. W skrócie program zajmuje się krzywizną obszarów: na każdy piksel obrazu wyjściowego narzucamy macierz, która określa kolor piksela wynikowego w generowanym obrazie z sąsiednich pikseli, używając mniej lub bardziej złożonej formuły. Zaletą tej funkcji jest to, że bardzo łatwo ją zrównoleglić, więc dany test doskonale pokazuje możliwości CUDA.

Kolejną zaletą jest to, że mamy już implementację na CPU, więc możemy porównać jej wynik z wersją CUDA - i nie wymyślać koła na nowo.

Powtarzamy raz jeszcze, że celem testu było zapoznanie się z narzędziami CUDA SDK, a nie porównanie wersji dla CPU i GPU. Ponieważ była to nasza pierwsza próba stworzenia programu CUDA, tak naprawdę nie spodziewaliśmy się wysokiej wydajności. Ponieważ ta część kodu nie jest krytyczna, wersja procesora nie została zoptymalizowana, więc bezpośrednie porównanie wyników jest mało interesujące.

Występ

Jednak zmierzyliśmy czas wykonania, aby zobaczyć, czy jest korzyść z używania CUDA nawet przy najtrudniejszej implementacji, czy też potrzebujemy długiej i żmudnej praktyki, aby uzyskać jakiekolwiek korzyści podczas korzystania z GPU. Maszyna testowa została zaczerpnięta z naszego laboratorium deweloperskiego - laptop z procesorem Core 2 Duo T5450 i kartą graficzną GeForce 8600M GT z systemem Vista. To daleko do superkomputera, ale wyniki są bardzo interesujące, ponieważ test nie jest „wyostrzony” dla GPU. Zawsze fajnie jest widzieć, jak nVidia pokazuje ogromne zyski w systemach z monstrualnymi procesorami graficznymi i dużą przepustowością, ale w praktyce wiele z 70 milionów procesorów graficznych obsługujących CUDA na dzisiejszym rynku komputerów PC nie jest tak potężnych, więc nasz test jest uzasadniony.

W przypadku obrazu o wymiarach 2048 x 2048 pikseli otrzymaliśmy następujące wyniki.

• Wątek procesora 1: 1419ms;
• Procesor 2 wątki: 749 ms;
• Procesor 4 wątki: 593ms
• GPU (8600M GT) bloki 256 wątków: 109ms;
• GPU (8600M GT) bloki 128 wątków: 94ms;
• Bloki GPU (8800 GTX) 128 wątków / 256 wątków: 31 ms.

Z uzyskanych wyników można wyciągnąć kilka wniosków. Zacznijmy od tego, że pomimo gadania o oczywistym lenistwie programistów, zmodyfikowaliśmy początkową wersję procesora dla kilku wątków. Jak już wspomnieliśmy, kod jest idealny do takiej sytuacji - wystarczy podzielić obraz początkowy na tyle stref, ile jest strumieni. Należy pamiętać, że od przejścia z jednego wątku do dwóch w naszym dwurdzeniowym procesorze przyspieszenie jest prawie liniowe, co również wskazuje na równoległy charakter programu testowego. Dość nieoczekiwanie wersja z czterema wątkami też była szybsza, choć jest to bardzo dziwne dla naszego procesora - wręcz przeciwnie, można było spodziewać się spadku wydajności ze względu na narzut związany z zarządzaniem dodatkowymi wątkami. Jak można wytłumaczyć taki wynik? Trudno powiedzieć, ale prawdopodobnie winowajcą jest harmonogram wątków Windows; w każdym razie wynik jest powtarzalny. Przy mniejszych teksturach (512x512) zysk z podziału nie był tak wyraźny (około 35% w porównaniu do 100%), a zachowanie wersji czterowątkowej było bardziej logiczne, bez wzmocnienia w porównaniu z wersją dwuwątkową. GPU był nadal szybszy, ale nie tak wyraźny (8600M GT był trzy razy szybszy niż wersja dwuwątkowa).

Drugą istotną obserwacją jest to, że nawet najwolniejsza implementacja GPU okazała się prawie sześciokrotnie szybsza niż najwydajniejsza wersja CPU. W przypadku pierwszego programu i niezoptymalizowanej wersji algorytmu wynik jest bardzo zachęcający. Należy pamiętać, że otrzymaliśmy namacalne najlepszy wynik na małych klockach, choć intuicja może podpowiadać inaczej. Wyjaśnienie jest proste - nasz program używa 14 rejestrów na wątek, a przy 256-wątkowych blokach wymagane jest 3584 rejestrów na blok, a 768 wątków jest wymaganych do pełnego obciążenia procesora, jak pokazaliśmy. W naszym przypadku są to trzy bloki lub 10 572 rejestry. Ale multiprocesor ma tylko 8192 rejestrów, więc może utrzymać aktywne tylko dwa bloki. W przeciwieństwie do bloków zawierających 128 wątków, potrzebujemy 1792 rejestrów na blok; jeśli 8192 podzielimy przez 1792 i zaokrąglimy w górę do najbliższej liczby całkowitej, otrzymamy cztery bloki. W praktyce liczba wątków będzie taka sama (512 na wieloprocesor, choć teoretycznie 768 jest potrzebne do pełnego obciążenia), ale zwiększenie liczby bloków daje GPU przewagę elastyczności w dostępie do pamięci - gdy operacja jest w toku z dużymi opóźnieniami może rozpocząć wykonywanie instrukcji innego bloku, czekając na otrzymanie wyników. Cztery bloki wyraźnie zmniejszają opóźnienia, zwłaszcza że nasz program korzysta z wielu dostępów do pamięci.

Analiza

Wreszcie, pomimo tego, co powiedzieliśmy powyżej, nie mogliśmy oprzeć się pokusie i uruchomiliśmy program na 8800 GTX, który był trzy razy szybszy niż 8600, niezależnie od rozmiaru bloku. Można by pomyśleć, że w praktyce na odpowiednich architekturach wynik będzie czterokrotnie lub więcej razy wyższy: 128 jednostek ALU/procesory cieniujące kontra 32 i wyższe taktowanie (1,35 GHz kontra 950 MHz), ale tak się nie stało. Najprawdopodobniej czynnikiem ograniczającym był dostęp do pamięci. Aby być bardziej precyzyjnym, początkowy obraz jest dostępny jako wielowymiarowa tablica CUDA - dość skomplikowany termin na coś, co jest niczym więcej niż teksturą. Ale jedzenie ma kilka korzyści.

• dostępy korzystają z pamięci podręcznej tekstur;
• używamy trybu zawijania, który nie musi obsługiwać granic obrazu, w przeciwieństwie do wersji procesora.

Ponadto możemy skorzystać z „darmowego” filtrowania ze znormalizowanym adresowaniem między zamiast i , ale w naszym przypadku jest to mało przydatne. Jak wiecie, 8600 ma 16 jednostek tekstur, w porównaniu do 32 dla 8800 GTX. Dlatego stosunek między dwiema architekturami wynosi tylko dwa do jednego. Dodajmy do tego różnicę częstotliwości i otrzymamy stosunek (32 x 0,575) / (16 x 0,475) = 2,4 - bliski rzeczywistemu "trzy do jednego". Ta teoria wyjaśnia również, dlaczego rozmiar bloków nie zmienia się zbytnio w G80, ponieważ ALU nadal opiera się na jednostkach tekstury.

Poza obiecującymi wynikami, nasze pierwsze spotkanie z CUDA przebiegło bardzo dobrze, biorąc pod uwagę wybrane niezbyt sprzyjające warunki. Programowanie na laptopie z systemem Vista oznacza korzystanie z CUDA SDK 2.0, który wciąż jest w wersji beta, ze sterownikiem 174.55, który również jest w fazie beta. Mimo to nie możemy zgłosić żadnych nieprzyjemnych niespodzianek - jedynie początkowe błędy podczas pierwszego debugowania, kiedy nasz program, wciąż bardzo "buggy", próbował adresować pamięć poza przydzielonym miejscem.

Monitor zaczął szaleńczo migotać, potem ekran pociemniał… aż Vista uruchomiła usługę naprawy sterowników i wszystko było w porządku. Ale nadal jest nieco zaskakujące, czy jesteś przyzwyczajony do typowego błędu segmentacji w standardowych programach, takich jak nasz. Na koniec mała krytyka pod adresem nVidii: w całej dokumentacji dostępnej dla CUDA nie ma małego przewodnika, który poprowadzi Cię krok po kroku, jak skonfigurować środowisko programistyczne dla Visual Studio. Właściwie problem jest niewielki, ponieważ SDK ma Pełen zestaw przykłady, które można zbadać, aby zrozumieć ramy dla aplikacji CUDA, ale przydałby się przewodnik dla początkujących.

Nvidia przedstawiła CUDA wraz z wydaniem GeForce 8800. I w tamtym czasie obietnice wydawały się bardzo kuszące, ale utrzymaliśmy nasz entuzjazm do prawdziwego testu. Rzeczywiście, w tamtym czasie pozostawanie na fali GPGPU wydawało się bardziej oznaczanie terytorium. Bez dostępnego pakietu SDK trudno powiedzieć, że nie mamy do czynienia z kolejnym marketingowym manekinem, który nie zadziała. To nie pierwszy raz, kiedy dobra inicjatywa została ogłoszona zbyt wcześnie i wtedy nie wyszła na jaw z powodu braku wsparcia – zwłaszcza w tak konkurencyjnym sektorze. Teraz, półtora roku po ogłoszeniu, możemy śmiało powiedzieć, że nVidia dotrzymała słowa.

SDK dość szybko wszedł do wersji beta na początku 2007 roku i od tego czasu był szybko aktualizowany, co dowodzi znaczenia tego projektu dla nVidii. Dziś CUDA rozwija się bardzo ładnie: SDK jest już dostępne w wersji beta 2.0 dla głównych systemów operacyjnych (Windows XP i Vista, Linux, a także 1.1 dla Mac OS X), a nVidia poświęciła całą sekcję witryny dla programiści.

Na bardziej profesjonalnym poziomie wrażenie pierwszych kroków z CUDA okazało się bardzo pozytywne. Nawet jeśli znasz architekturę GPU, możesz to łatwo rozgryźć. Kiedy interfejs API na pierwszy rzut oka wygląda przejrzyście, od razu zaczynasz wierzyć, że uzyskasz przekonujące wyniki. Ale czy czas obliczeniowy nie zostanie zmarnowany na wielokrotne transfery z CPU do GPU? A jak korzystać z tych tysięcy wątków prawie bez prymitywu synchronizacji? Rozpoczęliśmy nasze eksperymenty z myślą o tych wszystkich lękach. Ale szybko się rozproszyły, gdy pierwsza wersja naszego algorytmu, choć bardzo trywialna, okazała się znacznie szybsza niż na procesorze.

Tak więc CUDA nie jest ratunkiem dla naukowców, którzy chcą przekonać urzędników uniwersyteckich do zakupu GeForce. CUDA jest już w pełni dostępną technologią, z której może skorzystać każdy programista języka C, jeśli zechce poświęcić czas i wysiłek, aby przyzwyczaić się do nowego paradygmatu programowania. Te wysiłki nie pójdą na marne, jeśli twoje algorytmy będą dobrze zrównoleglone. Dziękujemy również firmie nVidia za dostarczenie pełnej i wysokiej jakości dokumentacji, w której początkujący programiści CUDA znajdą odpowiedzi.

Czego potrzeba, aby CUDA stała się rozpoznawalnym API? Jednym słowem: przenośność. Wiemy, że przyszłość IT leży w równoległym przetwarzaniu danych – dziś wszyscy już przygotowują się do takich zmian, a wszelkie inicjatywy, zarówno programowe, jak i sprzętowe, zmierzają w tym kierunku. Jednak w tej chwili, jeśli spojrzysz na rozwój paradygmatów, wciąż jesteśmy na początkowym etapie: ręcznie tworzymy wątki i staramy się planować dostęp do współdzielonych zasobów; z tym wszystkim można sobie jakoś poradzić, jeśli liczbę rdzeni można policzyć na palcach jednej ręki. Ale za kilka lat, kiedy liczba procesorów będzie liczyć w setkach, ta możliwość przestanie istnieć. Wraz z wydaniem CUDA, nVidia zrobiła pierwszy krok w rozwiązaniu tego problemu - ale oczywiście ta decyzja nadaje się tylko do GPU tej firmy, a nawet wtedy nie dla każdego. Tylko GF8 i 9 (oraz ich pochodne Quadro/Tesla) mogą dziś uruchamiać programy CUDA. I oczywiście nowa linia 260/280.

Nvidia może pochwalić się, że sprzedała na całym świecie 70 milionów procesorów graficznych zgodnych z CUDA, ale to wciąż za mało, aby stać się de facto standardem. Biorąc pod uwagę fakt, że konkurenci nie siedzą bezczynnie. AMD oferuje własne SDK (Stream Computing), a Intel ogłosił rozwiązanie (Ct), chociaż nie jest ono jeszcze dostępne. Zbliża się wojna o standardy i wyraźnie nie będzie miejsca na rynku dla trzech konkurentów, dopóki inny gracz, taki jak Microsoft, nie przedstawi wspólnej propozycji API, co oczywiście ułatwi życie programistom.

Dlatego nVidia ma wiele trudności z zatwierdzeniem przez CUDA. Choć technologicznie mamy przed sobą bez wątpienia udane rozwiązanie, to pozostaje przekonać deweloperów o jego perspektywach – a to nie będzie łatwe. Jednak sądząc po wielu ostatnich ogłoszeniach i nowościach API, przyszłość wcale nie wygląda ponuro.

Technologia CUDA

Władimir Frołow,[e-mail chroniony]

adnotacja

Artykuł mówi o technologii CUDA, która pozwala programiście używać kart graficznych jako potężnych jednostek obliczeniowych. Narzędzia dostarczane przez Nvidię umożliwiają pisanie programów procesorów graficznych (GPU) w podzbiorze języka C++. To zwalnia programistę z konieczności używania shaderów i zrozumienia działania potoku graficznego. W artykule przedstawiono przykłady programowania z wykorzystaniem CUDA oraz różne techniki optymalizacji.

1. Wstęp

Rozwój technologii komputerowych w ciągu ostatnich dziesięcioleci przebiegał w szybkim tempie. Tak szybko, że twórcy procesorów niemal doszli do tak zwanego „krzemowego ślepego zaułka”. Nieokiełznany wzrost częstotliwości zegara stał się niemożliwy z wielu poważnych przyczyn technologicznych.

Częściowo dlatego wszyscy producenci nowoczesnych systemów komputerowych dążą do zwiększania liczby procesorów i rdzeni, zamiast zwiększania częstotliwości jednego procesora. Liczba rdzeni jednostki centralnej (CPU) w zaawansowanych systemach wynosi już 8.

Innym powodem jest stosunkowo niska prędkość pamięci RAM. Bez względu na to, jak szybko działa procesor, wąskimi gardłami, jak pokazuje praktyka, nie są wcale operacje arytmetyczne, ale nieudane dostępy do pamięci - braki w pamięci podręcznej.

Jeśli jednak spojrzysz w kierunku GPU (Graphics Processing Unit), to znacznie wcześniej poszli ścieżką równoległości. Dzisiejsze karty graficzne, takie jak GF8800GTX, mogą mieć procesory do 128. Wydajność takich systemów, jeśli zostanie umiejętnie zaprogramowana, może być dość znacząca (rys. 1).

Ryż. 1. Liczba operacji zmiennoprzecinkowych dla CPU i GPU

Kiedy na rynku pojawiły się pierwsze karty graficzne, były one dość prostymi (w porównaniu z procesorem centralnym), wysoce wyspecjalizowanymi urządzeniami zaprojektowanymi do przejęcia ciężaru wizualizacji dwuwymiarowych danych z procesora. Wraz z rozwojem branży gier i pojawieniem się takich trójwymiarowych gier, jak Doom (rys. 2) i Wolfenstein 3D (rys. 3), pojawiła się potrzeba wizualizacji 3D.

Rysunki 2.3. Gry 3D Doom i Wolfenstein

Od czasu stworzenia pierwszych kart graficznych Voodoo przez 3Dfx (1996) i do 2001 roku w GPU zaimplementowano tylko stały zestaw operacji na danych wejściowych.

Programiści nie mieli wyboru w algorytmie renderowania, a dla zwiększenia elastyczności pojawiły się shadery - małe programy, które są wykonywane przez kartę graficzną dla każdego wierzchołka lub dla każdego piksela. Do ich zadań należały przekształcenia wierzchołków i cieniowanie - obliczanie oświetlenia w punkcie np. według modelu Phonga.

Chociaż shadery są obecnie bardzo zaawansowane, należy rozumieć, że zostały opracowane z myślą o wysoce wyspecjalizowanych zadaniach związanych z transformacją i rasteryzacją 3D. Podczas gdy procesory graficzne ewoluują w kierunku uniwersalnych systemów wieloprocesorowych, języki shaderów pozostają wysoce wyspecjalizowane.

Można je porównać do języka FORTRAN w tym sensie, że podobnie jak FORTRAN, były pierwsze, ale przeznaczone do rozwiązywania tylko jednego rodzaju problemu. Shadery są mało przydatne do rozwiązywania problemów innych niż transformacje 3D i rasteryzacja, podobnie jak FORTRAN nie nadaje się do rozwiązywania problemów niezwiązanych z obliczeniami numerycznymi.

Obecnie istnieje trend nietradycyjnego wykorzystywania kart wideo do rozwiązywania problemów z zakresu mechaniki kwantowej, sztucznej inteligencji, obliczeń fizycznych, kryptografii, fizycznie poprawnej wizualizacji, rekonstrukcji ze zdjęć, rozpoznawania itp. Zadania te są niewygodne do rozwiązania w graficznych API (DirectX, OpenGL), ponieważ te API zostały stworzone dla zupełnie innych aplikacji.

Rozwój programowania ogólnego na GPU (GPGPU) logicznie doprowadził do pojawienia się technologii mających na celu szerszy zakres zadań niż rasteryzacja. W rezultacie Nvidia stworzyła technologię Compute Unified Device Architecture (w skrócie CUDA), a konkurencyjna firma ATI stworzyła technologię STREAM.

Należy zauważyć, że w momencie pisania tego artykułu technologia STREAM była daleko w rozwoju w rozwoju CUDA i dlatego nie będzie tutaj brana pod uwagę. Skoncentrujemy się na CUDA, technologii GPGPU, która umożliwia pisanie programów w podzbiorze języka C++.

2. Podstawowa różnica między CPU a GPU

Rozważmy pokrótce niektóre znaczące różnice między obszarami i funkcjami aplikacji procesora centralnego i karty graficznej.

2.1. Możliwości

Procesor jest oryginalnie przystosowany do rozwiązywania ogólnych zadań i współpracuje z dowolnie adresowalną pamięcią. Programy na procesorze mają bezpośredni dostęp do dowolnych komórek pamięci liniowej i jednorodnej.

W przypadku GPU tak nie jest. Jak przekonasz się po przeczytaniu tego artykułu, w CUDA występuje aż 6 rodzajów pamięci. Możesz czytać z dowolnej komórki, która jest fizycznie dostępna, ale nie możesz pisać do wszystkich komórek. Powodem jest to, że GPU jest w każdym razie konkretnym urządzeniem zaprojektowanym do określonych celów. To ograniczenie zostało wprowadzone w celu zwiększenia szybkości niektórych algorytmów i obniżenia kosztów sprzętu.

2.2. Wydajność pamięci

Odwiecznym problemem większości systemów komputerowych jest to, że pamięć jest wolniejsza niż procesor. Producenci procesorów rozwiązują go, wprowadzając pamięci podręczne. Najczęściej używane obszary pamięci są umieszczane w pamięci podręcznej lub pamięci podręcznej, działającej z częstotliwością procesora. Pozwala to zaoszczędzić czas przy dostępie do najczęściej używanych danych i załadować do procesora rzeczywiste obliczenia.

Zauważ, że pamięci podręczne są w rzeczywistości niewidoczne dla programisty. Zarówno podczas odczytu, jak i zapisu dane nie trafiają od razu do pamięci RAM, ale przechodzą przez pamięci podręczne. Pozwala to w szczególności na szybkie odczytanie jakiejś wartości zaraz po zapisaniu.

Na GPU (tu mamy na myśli karty graficzne GF z serii ósmej) są też pamięci podręczne i one również są ważne, ale ten mechanizm nie jest tak potężny, jak na procesorze. Po pierwsze, nie wszystkie typy pamięci są buforowane, a po drugie, pamięci podręczne są tylko do odczytu.

Na GPU powolne dostępy do pamięci są ukryte przy użyciu obliczeń równoległych. Podczas gdy niektóre zadania czekają na dane, inne działają, gotowe do obliczeń. Jest to jedna z głównych zasad CUDA, która pozwala znacznie poprawić wydajność systemu jako całości.

3. rdzeń CUDA

3.1. model przesyłania strumieniowego

Architektura obliczeniowa CUDA opiera się na koncepcjijedno polecenie dla dużej ilości danych(Single Instruction Multiple Data, SIMD) i koncepcja wieloprocesorowy.

Koncepcja SIMD zakłada, że ​​jedna instrukcja pozwala na przetwarzanie wielu danych jednocześnie. Na przykład polecenie addps w Pentium 3 i nowszych Pentiumach umożliwia jednoczesne dodanie 4 liczb zmiennoprzecinkowych pojedynczej precyzji.

Multiprocesor to wielordzeniowy procesor SIMD, który pozwala na wykonanie tylko jednej instrukcji na wszystkich rdzeniach w danym momencie. Każdy rdzeń wieloprocesorowy jest skalarny, tj. nie obsługuje operacji wektorowych w najczystszej postaci.

Zanim przejdziemy dalej, przedstawmy kilka definicji. Zwróć uwagę, że urządzenie i host w tym artykule nie zostaną w ogóle zrozumiane, jak większość programistów jest do tego przyzwyczajona. Będziemy używać takich terminów, aby uniknąć rozbieżności z dokumentacją CUDA.

Pod urządzeniem (urządzeniem) w naszym artykule zrozumiemy kartę wideo obsługującą sterownik CUDA lub inne wyspecjalizowane urządzenie przeznaczone do wykonywania programów za pomocą CUDA (takie jak na przykład NVIDIA Tesla). W naszym artykule rozważymy GPU tylko jako urządzenie logiczne, unikając konkretnych szczegółów implementacji.

Host (host) nazwiemy programem w zwykłej pamięci RAM komputera, który wykorzystuje procesor i wykonuje funkcje sterujące do pracy z urządzeniem.

W rzeczywistości część twojego programu działająca na procesorze to gospodarz, i twoja karta wideo urządzenie. Logicznie rzecz biorąc, urządzenie można traktować jako zestaw multiprocesorów (rysunek 4) plus sterownik CUDA.

Ryż. 4. Urządzenie

Załóżmy, że chcemy uruchomić określoną procedurę na naszym urządzeniu w N wątkach (czyli chcemy zrównoleglić jej pracę). Zgodnie z dokumentacją CUDA, nazwijmy tę procedurę jądrem.

Cechą architektury CUDA jest organizacja blokowo-siatkowa, co jest nietypowe dla aplikacji wielowątkowych (rys. 5). Jednocześnie sterownik CUDA niezależnie rozdziela zasoby urządzeń między wątki.

Ryż. 5. Organizacja przepływów

Na ryc. 5. rdzeń jest oznaczony jako Kernel. Wszystkie wątki realizujące ten rdzeń są łączone w bloki (Blok), a bloki z kolei są łączone w siatkę (Siatka).

Jak widać na rysunku 5, do identyfikacji przepływów wykorzystywane są wskaźniki dwuwymiarowe. Twórcy CUDA udostępnili możliwość pracy z indeksami trójwymiarowymi, dwuwymiarowymi lub prostymi (jednowymiarowymi) w zależności od tego, co jest wygodniejsze dla programisty.

Ogólnie rzecz biorąc, indeksy są trójwymiarowymi wektorami. Dla każdego wątku będą znane: indeks wątku w bloku threadIdx i indeks bloku w siatce blockIdx. Na starcie wszystkie wątki będą się różnić tylko tymi indeksami. W rzeczywistości to za pomocą tych indeksów programista sprawuje kontrolę, określając, która część jego danych jest przetwarzana w każdym wątku.

Odpowiedź na pytanie, dlaczego twórcy wybrali taką organizację, nie jest banalna. Jednym z powodów jest gwarantowane wykonanie jednego bloku. na jednego urządzenie wieloprocesorowe, ale jeden multiprocesor może wykonywać kilka różnych bloków. Inne powody zostaną wyjaśnione w dalszej części artykułu.

Blok zadań (wątków) jest wykonywany na wieloprocesorze w częściach lub pulach, zwanych wypaczeniami. Obecny rozmiar wypaczenia w kartach graficznych z obsługą CUDA to 32 wątki. Zadania wewnątrz puli warp są wykonywane w stylu SIMD, tj. wszystkie wątki wewnątrz wypaczenia mogą wykonywać tylko jedną instrukcję na raz.

Tutaj należy zrobić jedno zastrzeżenie. W architekturach nowoczesnych w momencie pisania tego tekstu liczba procesorów w jednym multiprocesorze wynosi 8, a nie 32. Wynika z tego, że nie całe osnowa jest wykonywana w tym samym czasie, jest podzielona na 4 części, które są wykonywane sekwencyjnie (ponieważ procesory są skalarne) .

Ale po pierwsze, twórcy CUDA nie regulują ściśle rozmiaru wypaczenia. W swoich pracach wspominają o parametrze warp size, a nie o liczbie 32. Po drugie, z logicznego punktu widzenia jest to osnowa, czyli minimalna suma wątków, o której można powiedzieć, że wszystkie wątki w niej zawarte są wykonywane jednocześnie - a jednocześnie nie przyjmuje się żadnych założeń, że system odpoczynku nie zostanie wykonany.

3.1.1. rozgałęzienia

Natychmiast pojawia się pytanie: jeśli w tym samym czasie wszystkie wątki wewnątrz osnowy wykonują tę samą instrukcję, to co z gałęziami? W końcu, jeśli w kodzie programu wystąpi rozgałęzienie, instrukcje będą już inne. Tutaj zastosowano standardowe rozwiązanie do programowania SIMD (rys. 6).

Ryż. 6. Organizacja rozgałęzień w SIMD

Miejmy następujący kod:

jeśli (warunki)B;

W przypadku SISD (Single Instruction Single Data) wykonujemy instrukcję A, sprawdzamy warunek, a następnie wykonujemy instrukcje B i D (jeśli warunek jest spełniony).

Załóżmy teraz, że mamy 10 wątków działających w stylu SIMD. We wszystkich 10 wątkach wykonujemy instrukcję A, następnie sprawdzamy warunek cond i okazuje się, że w 9 na 10 wątków jest to prawda, aw jednym fałsz.

Jasne jest, że nie możemy uruchomić 9 wątków w celu wykonania instrukcji B i jednego pozostałego wątku w celu wykonania instrukcji C, ponieważ tylko jedna instrukcja może być wykonana jednocześnie we wszystkich wątkach. W takim przypadku musisz to zrobić: najpierw „zabijamy” podzielony wątek, aby nie zepsuł niczyich danych, i wykonujemy 9 pozostałych wątków. Następnie „zabijamy” 9 wątków, które wykonały instrukcję B, i przechodzimy przez jeden wątek z instrukcją C. Następnie wątki są ponownie łączone i wykonują jednocześnie instrukcję D.

Okazuje się, że jest to smutny wynik: zasoby procesora są nie tylko zużywane na szlifowanie pustych bitów w dzielonych strumieniach, ale, co gorsza, w końcu będziemy zmuszeni do wykonania OBU gałęzi.

Jednak nie wszystko jest tak złe, jak mogłoby się wydawać na pierwszy rzut oka. Bardzo dużą zaletą technologii jest to, że triki te wykonywane są dynamicznie przez sterownik CUDA i są całkowicie transparentne dla programisty. Jednocześnie, mając do czynienia z instrukcjami SSE nowoczesnych procesorów (czyli w przypadku próby wykonania 4 kopii algorytmu jednocześnie) sam programista musi zadbać o szczegóły: łączyć dane czterokrotnie, robić nie zapomnij o wyrównaniu i generalnie pisz na niskim poziomie w zasadzie jak w asemblerze.

Z powyższego wynika jeden bardzo ważny wniosek. Gałęzie same w sobie nie powodują obniżenia wydajności. Szkodliwe są tylko gałęzie, w których wątki rozchodzą się w ramach tej samej puli wątków osnowy. Co więcej, jeśli wątki rozchodzą się w tym samym bloku, ale w różnych pulach osnowy lub w różnych blokach, nie ma to absolutnie żadnego efektu.

3.1.2. Komunikacja między wątkami

W momencie pisania tego artykułu jakakolwiek interakcja między wątkami (synchronizacja i wymiana danych) była możliwa tylko w obrębie bloku. Oznacza to, że niemożliwe jest zorganizowanie interakcji między przepływami różnych bloków przy użyciu wyłącznie udokumentowanych funkcji.

Jeśli chodzi o funkcje nieudokumentowane, zdecydowanie odradza się ich używanie. Powodem tego jest to, że opierają się one na specyficznych cechach sprzętowych konkretnego systemu.

Synchronizacja wszystkich zadań w bloku odbywa się poprzez wywołanie funkcji __synchtreads. Wymiana danych jest możliwa dzięki pamięci współdzielonej, ponieważ jest wspólna dla wszystkich zadań w obrębie bloku.

3.2. Pamięć

W CUDA istnieje sześć typów pamięci (rys. 7). Są to rejestry, pamięć lokalna, globalna, współdzielona, ​​stała i tekstura.

Taka obfitość wynika ze specyfiki karty graficznej i jej głównego celu, a także z chęci programistów, aby system był jak najtańszy, poświęcając w różnych przypadkach wszechstronność lub szybkość.

Ryż. 7. Rodzaje pamięci w CUDA

3.2.0. Rejestry

Jeśli to możliwe, kompilator próbuje umieścić wszystkie lokalne zmienne funkcji w rejestrach. Dostęp do takich zmiennych odbywa się z maksymalną prędkością. W obecnej architekturze na jeden procesor przypada 8192 rejestrów 32-bitowych. Aby określić ile rejestrów jest dostępnych dla jednego wątku należy tę liczbę (8192) podzielić przez rozmiar bloku (liczbę wątków w nim zawartych).

Przy typowym podziale 64 wątków na blok, jest tylko 128 rejestrów (istnieją pewne obiektywne kryteria, ale 64 jest odpowiednie średnio dla wielu zadań). Realistycznie rzecz biorąc, nvcc nigdy nie przydzieli 128 rejestrów. Zwykle nie daje więcej niż 40, a reszta zmiennych trafi do pamięci lokalnej. Dzieje się tak, ponieważ na jednym multiprocesorze można wykonać kilka bloków. Kompilator próbuje zmaksymalizować liczbę jednocześnie działających bloków. Dla większej wydajności należy starać się zajmować mniej niż 32 rejestry. Wtedy teoretycznie na jednym multiprocesorze można uruchomić 4 bloki (8 wypaczeń, jeśli 64 wątki w jednym bloku). Jednak należy również wziąć pod uwagę ilość pamięci współdzielonej zajmowanej przez wątki, ponieważ jeśli jeden blok zajmuje całą pamięć współdzieloną, dwa takie bloki nie mogą być wykonywane jednocześnie na wieloprocesorze.

3.2.1. pamięć lokalna

W przypadkach, gdy lokalne dane procedur są zbyt duże lub kompilator nie może obliczyć dla nich jakiegoś stałego kroku podczas uzyskiwania do nich dostępu, może umieścić je w pamięci lokalnej. Można to ułatwić, na przykład, rzucając wskaźniki na typy o różnych rozmiarach.

Fizycznie pamięć lokalna jest analogiczna do pamięci globalnej i działa z tą samą prędkością. W chwili pisania tego tekstu nie istniały żadne mechanizmy, które wyraźnie zakazywałyby kompilatorowi używania pamięci lokalnej dla określonych zmiennych. Ponieważ kontrola pamięci lokalnej jest dość trudna, lepiej w ogóle jej nie używać (patrz Rozdział 4 „Zalecenia dotyczące optymalizacji”).

3.2.2. pamięć globalna

W dokumentacji CUDA jako jedno z głównych osiągnięćTechnologia zapewnia możliwość dowolnego adresowania pamięci globalnej. Oznacza to, że możesz czytać z dowolnej komórki pamięci, a także możesz pisać do dowolnej komórki (zwykle nie ma to miejsca na GPU).

Jednak za wszechstronność w tym przypadku trzeba płacić szybkością. Pamięć globalna nie jest buforowana. Działa bardzo wolno, mimo to liczba globalnych dostępów do pamięci powinna być zminimalizowana.

Pamięć globalna jest potrzebna głównie do zapisywania wyników programu przed wysłaniem ich do hosta (do zwykłej pamięci DRAM). Powodem tego jest to, że pamięć globalna jest jedynym rodzajem pamięci, w której można coś zapisać.

Zmienne zadeklarowane z kwalifikatorem __global__ są umieszczane w pamięci globalnej. Pamięć globalną można również alokować dynamicznie, wywołując funkcję cudaMalloc(void* mem, int size) na hoście. Nie można wywołać tej funkcji z urządzenia. Wynika z tego, że alokacja pamięci powinna być obsługiwana przez program hosta działający na procesorze. Dane z hosta można przesłać do urządzenia poprzez wywołanie funkcji cudaMemcpy:

cudaMemcpy(void* gpu_mem, void* cpu_mem, int rozmiar, cudaMemcpyHostToDevice);

W ten sam sposób możesz wykonać odwrotną procedurę:

cudaMemcpy(void* cpu_mem, void* gpu_mem, int rozmiar, cudaMemcpyDeviceToHost);

To połączenie jest również wykonywane z hosta.

Podczas pracy z pamięcią globalną ważne jest przestrzeganie zasady koalescencji. Główną ideą jest to, że wątki powinny mieć dostęp do kolejnych komórek pamięci i 4,8 lub 16 bajtów. W takim przypadku pierwszy wątek musi adresować adres wyrównany do granicy, odpowiednio 4,8 lub 16 bajtów. Adresy zwracane przez cudaMalloc są wyrównane do granicy co najmniej 256 bajtów.

3.2.3. Pamięć współdzielona

Pamięć współdzielona jest pamięcią niebuforowaną, ale szybką. Zaleca się używanie go jako zarządzanej pamięci podręcznej. Tylko 16 KB pamięci współdzielonej jest dostępne na wieloprocesor. Dzieląc tę ​​liczbę przez liczbę zadań w bloku otrzymujemy maksymalna ilość pamięć współdzielona dostępna na wątek (jeśli planujesz używać jej niezależnie we wszystkich wątkach).

Cechą charakterystyczną pamięci współdzielonej jest to, że adresowana jest tak samo dla wszystkich zadań w obrębie bloku (Rysunek 7). Wynika z tego, że może służyć do wymiany danych między wątkami tylko jednego bloku.

Gwarantuje to, że podczas wykonywania bloku na wieloprocesorze zostanie zachowana zawartość pamięci współdzielonej. Jednak po zmianie bloku na wieloprocesorze nie ma gwarancji, że zawartość starego bloku zostanie zachowana. Dlatego nie należy próbować synchronizować zadań między blokami, pozostawiając jakiekolwiek dane w pamięci współdzielonej i licząc na ich bezpieczeństwo.

Zmienne zadeklarowane z kwalifikatorem __shared__ są umieszczane w pamięci współdzielonej.

shared_float mem_shared;

Należy jeszcze raz podkreślić, że na blok przypada tylko jedna pamięć współdzielona. Dlatego jeśli chcesz używać go tylko jako zarządzanej pamięci podręcznej, powinieneś odwoływać się do różnych elementów tablicy, na przykład tak:

pływak x = współdzielona_pamięć;

Gdzie threadIdx.x jest indeksem x wątku wewnątrz bloku.

3.2.4. Stała pamięć

Pamięć stała jest buforowana, jak pokazano na ryc. 4. Pamięć podręczna istnieje w jednej instancji dla jednego wieloprocesora, co oznacza, że ​​jest wspólna dla wszystkich zadań w obrębie bloku. Na hoście możesz zapisać coś do pamięci stałej, wywołując funkcję cudaMemcpyToSymbol. Z urządzenia pamięć stała jest tylko do odczytu.

Pamięć stała jest bardzo wygodna w użyciu. Możesz umieścić w nim dane dowolnego typu i odczytać je za pomocą prostego przypisania.

#zdefiniuj N 100

Constant__int gpu_buffer[N];

nieważna funkcja_hosta()

int bufor_cpu[N];

cudaMemcpyToSymbol(gpu_buffer, cpu_buffer, sizeof(int )*N);

// __global__ oznacza, że ​​device_kernel to jądro, które można uruchomić na GPU

Global__void device_kernel()

int a = bufor_gpu;

int b = gpu_buffer + gpu_buffer;

// gpu_buffer = a; BŁĄD! pamięć stała jest tylko do odczytu

Ponieważ pamięć podręczna jest używana jako pamięć stała, dostęp do niej jest zazwyczaj dość szybki. Jedyną, ale bardzo dużą wadą pamięci stałej jest to, że jej rozmiar to tylko 64 KB (dla całego urządzenia). Wynika z tego, że sensowne jest przechowywanie tylko niewielkiej ilości często używanych danych w pamięci kontekstowej.

3.2.5. pamięć tekstur

Pamięć tekstur jest buforowana (rys. 4). Na każdy procesor wieloprocesorowy przypada tylko jedna pamięć podręczna, co oznacza, że ​​ta pamięć podręczna jest współdzielona przez wszystkie zadania w bloku.

Nazwa pamięci tekstur (i niestety funkcjonalności) jest dziedziczona z pojęć „tekstura” i „teksturowanie”. Teksturowanie to proces nakładania tekstury (tylko obrazu) na wielokąt podczas rasteryzacji. Pamięć tekstur jest zoptymalizowana pod kątem próbkowania danych 2D i ma następujące funkcje:

szybki wybór wartości o stałym rozmiarze (bajt, słowo, słowo podwójne lub poczwórne) z tablicy jednowymiarowej lub dwuwymiarowej;

znormalizowane adresowanie z liczbami zmiennoprzecinkowymi w przedziale . Możesz je następnie wybrać za pomocą znormalizowanego adresowania. Wynikowa wartość będzie słowem float4 odwzorowanym na interwał;

CudaMalloc((void**) &gpu_memory, N*sizeof (uint4 )); // przydziel pamięć w GPU

// ustawienie parametrów tekstury tekstury

Texture.addressMode = cudaAddressModeWrap; // tryb Zawinąć

Texture.addressMode = cudaAddressModeWrap;

Texture.filterMode = cudaFilterModePoint; //najbliższa wartość

texture.normalized = fałsz; // nie używaj znormalizowanego adresowania

CudaBindTexture(0, tekstura , gpu_memory , N ) // odtąd ta pamięć będzie uważana za pamięć tekstur


cudaMemcpy(gpu_memory, cpu_buffer, N*sizeof(uint 4), cudaMemcpyHostToDevice ); // skopiuj dane doGPU

// __global__ oznacza, że ​​device_kernel jest jądrem do zrównoleglenia

Global__void device_kernel()

uint4 a = tex1Dpobierz(tekstura,0); // możesz wybrać dane tylko w ten sposób!

uint4 b = tex1Dpobierz(tekstura;1);

int c = a.x * b.y;

...

3.3. Prosty przykład

Jako prosty przykład rozważmy program cppIntegration z pakietu CUDA SDK. Pokazuje, jak pracować z CUDA, a także używać nvcc (specjalny kompilator podzbiorów C++ firmy NVIDIA) w połączeniu z MS Visual Studio, co znacznie upraszcza tworzenie programów CUDA.

4.1. Podziel poprawnie swoje zadanie

Nie wszystkie zadania są odpowiednie dla architektur SIMD. Jeśli twoje zadanie nie nadaje się do tego, użycie GPU może nie być opłacalne. Ale jeśli jesteś zdecydowany używać GPU, powinieneś spróbować podzielić algorytm na części, które można wydajnie wykonać w stylu SIMD. Jeśli to konieczne, zmień algorytm, aby rozwiązać swój problem, wymyśl nowy - taki, który dobrze pasowałby do SIMD. Przykładem odpowiedniego przypadku użycia dla GPU jest implementacja piramidowego dodawania elementów tablicy.

4.2. Wybór typu pamięci

Umieść swoje dane w teksturach lub pamięci stałej, jeśli wszystkie zadania w tym samym bloku uzyskują dostęp do tej samej lokalizacji pamięci lub zamknij lokalizacje. Dane dwuwymiarowe można efektywnie przetwarzać za pomocą funkcji text2Dfetch i text2D. Pamięć tekstur jest specjalnie zoptymalizowana pod kątem próbkowania 2D.

Użyj pamięci globalnej w połączeniu z pamięcią współdzieloną, jeśli wszystkie zadania losowo uzyskują dostęp do różnych, szeroko rozstawionych obszarów pamięci (z bardzo różnymi adresami lub współrzędnymi, jeśli są to dane 2D/3D).

pamięć globalna => pamięć współdzielona

syncthreads();

Przetwarzaj dane w pamięci współdzielonej

syncthreads();

pamięć globalna<= разделяемая память

4.3. Włącz liczniki pamięci

Flaga kompilatora --ptxas-options=-v pozwala dokładnie określić, ile i jakiego rodzaju pamięci (rejestry, współdzielona, ​​lokalna, stała) używasz. Jeśli kompilator korzysta z pamięci lokalnej, wiesz dokładnie, co to jest. Analiza danych o ilości i typach używanej pamięci może znacznie pomóc w optymalizacji programu.

4.4. Postaraj się zminimalizować użycie rejestrów i pamięci współdzielonej

Im bardziej jądro używa rejestrów lub pamięci dzielonej, tym mniej wątków (a raczej wypaczeń) może jednocześnie działać na wieloprocesorowym, ponieważ. zasoby wieloprocesorowe są ograniczone. Dlatego niewielki wzrost zajętości rejestrów lub pamięci współdzielonej może w niektórych przypadkach doprowadzić do spadku wydajności o połowę - właśnie dlatego, że obecnie na wieloprocesorze jest jednocześnie wykonywanych dokładnie dwa razy mniej wypaczeń.

4.5. Pamięć współdzielona zamiast lokalnej

Jeśli kompilator Nvidii z jakiegoś powodu zaalokował dane w pamięci lokalnej (zazwyczaj jest to zauważalne bardzo silnym spadkiem wydajności w miejscach, w których nie ma nic zasobochłonnego), dowiedz się dokładnie, jakie dane dostały się do pamięci lokalnej i umieść je we współdzielonej pamięć (pamięć współdzielona).

Często kompilator alokuje zmienną w pamięci lokalnej, jeśli nie jest ona często używana. Na przykład jest to rodzaj akumulatora, w którym akumulujesz wartość, obliczając coś w pętli. Jeśli pętla ma duży rozmiar kodu (ale nie czas wykonania!), kompilator może umieścić Twój akumulator w pamięci lokalnej, ponieważ jest stosunkowo rzadko używany, a rejestrów jest niewiele. Utrata wydajności w tym przypadku może być zauważalna.

Jeśli naprawdę rzadko używasz zmiennej, lepiej umieścić ją jawnie w pamięci globalnej.

Chociaż może się wydawać, że kompilator automatycznie alokuje takie zmienne w pamięci lokalnej, w rzeczywistości tak nie jest. Nie będzie łatwo znaleźć wąskie gardło przy kolejnych modyfikacjach programu, jeśli zmienna zacznie być częściej używana. Kompilator może, ale nie musi, przenieść taką zmienną do pamięci rejestru. Jeśli modyfikator __global__ jest określony wprost, programista z większym prawdopodobieństwem zwróci na niego uwagę.

4.6. Rozwijanie pętli

Rozwijanie pętli to standardowa sztuczka wydajnościowa w wielu systemach. Jego istotą jest wykonanie większej liczby akcji w każdej iteracji, zmniejszając w ten sposób całkowitą liczbę iteracji, a co za tym idzie liczbę skoków warunkowych, które procesor będzie musiał wykonać.

Oto jak możesz rozwinąć pętlę w celu znalezienia sumy tablicy (na przykład liczby całkowitej):

int a[N]; insum;

dla (int i=0;i

Oczywiście pętle można rozwijać ręcznie (jak pokazano powyżej), ale jest to bezproduktywna praca. O wiele lepiej jest używać szablonów C++ w połączeniu z funkcjami wbudowanymi.

szablon

class ArraySuma

Device__ static T exec(const T* arr) ( return arr + ArraySumm (arr+1); )

szablon

class ArraySuma<0,T>

Device__ static T exec(const T* arr) ( zwraca 0; )

dla (int i=0;i

suma+= ArraySuma<4,int>::exec(a);

Należy zwrócić uwagę na jedną interesującą cechę kompilatora nvcc. Kompilator zawsze będzie domyślnie wbudowany w funkcje, takie jak __device__ (jest specjalna dyrektywa __noinline__, która to przesłoni).

Dlatego możesz być pewien, że przykład taki jak ten powyżej rozwinie się w prostą sekwencję instrukcji i nie będzie w żaden sposób gorszy od ręcznie napisanego kodu. Jednak w ogólnym przypadku (nie nvcc) nie możesz być tego pewien, ponieważ inline jest tylko wskazówką dla kompilatora, którą może zignorować. Dlatego nie ma gwarancji, że Twoje funkcje będą wbudowane.

4.7. Wyrównanie danych i próbkowanie 16-bajtowe

Wyrównaj struktury danych na granicy 16 bajtów. W takim przypadku kompilator będzie mógł użyć dla nich specjalnych instrukcji, które ładują dane od razu w 16 bajtach.

Jeśli struktura ma 8 bajtów lub mniej, możesz wyrównać ją do 8 bajtów. Ale w tym przypadku możesz wybrać dwie zmienne naraz, łącząc dwie 8-bajtowe zmienne w strukturę za pomocą rzutowania unii lub wskaźnika. Rzutowanie powinno być używane ostrożnie, ponieważ kompilator może umieszczać dane w pamięci lokalnej, a nie w rejestrach.

4.8. Konflikty banków pamięci współdzielonej

Pamięć współdzielona jest zorganizowana w postaci 16 (tylko!) banków pamięci z krokiem 4 bajtów. Podczas wykonywania puli wątków warp na wieloprocesorze jest ona dzielona na dwie połowy (jeśli warp-size = 32) z 16 wątków, które kolejno uzyskują dostęp do pamięci współdzielonej.

Zadania w różnych połówkach osnowy nie powodują konfliktu w pamięci współdzielonej. Ze względu na to, że zadania połowy puli warp będą miały dostęp do tych samych banków pamięci, wystąpią kolizje i w rezultacie spadnie wydajność. Zadania w tej samej połowie warp mogą uzyskiwać dostęp do różnych obszarów pamięci współdzielonej z pewnym krokiem.

Optymalne kroki to 4, 12, 28, ..., 2^n-4 bajty (rys. 8).



Ryż. 8. Optymalne kroki.

Nieoptymalne kroki to 1, 8, 16, 32, ..., 2^n bajtów (rys. 9).



Ryż. 9. Nieoptymalne kroki

4.9. Minimalizowanie przesunięć danych hosta<=>urządzenie

Staraj się przesyłać jak najmniej wyników pośrednich do hosta w celu przetworzenia przez procesor. Implementuj, jeśli nie cały algorytm, to przynajmniej jego główną część na GPU, pozostawiając CPU tylko zadania kontrolne.

5. Przenośna biblioteka matematyczna CPU/GPU

Autor tego artykułu napisał przenośną bibliotekę MGML_MATH do pracy z prostymi obiektami przestrzennymi, której kod działa zarówno na urządzeniu, jak i na hoście.

Biblioteka MGML_MATH może być używana jako platforma do pisania przenośnych (lub hybrydowych) systemów CPU/GPU do obliczania fizycznych, graficznych lub innych problemów przestrzennych. Jego główną zaletą jest to, że ten sam kod może być używany zarówno na CPU, jak i na GPU, a jednocześnie na czele wymagań stawianych bibliotece stoi szybkość.

6 . Literatura

Chrisa Kaspersky’ego. Technika optymalizacji programu. Efektywne wykorzystanie pamięci. - St. Petersburg: BHV-Petersburg, 2003. - 464 s.: ch.

Przewodnik programowania CUDA 1.1 ( http://developer.download.nvidia.com/compute/cuda/1_1/NVIDIA_CUDA_Programming_Guide_1.1.pdf )

Przewodnik programowania CUDA 1.1. strony 14-15

Przewodnik programowania CUDA 1.1. strona 48