tanulas_cuda_programozasi_alapismeretek
Fehér Krisztián: CUDA programozási alapismeretek
Tartalomjegyzék
1 Előszó
1.1 Célközönség, követelmények
1.2 Letölthető programkódok
2 A CUDA platform bemutatása
2.1 Mi az a CUDA?
2.2 CUDA alapú eszközök
2.3 CUDA eszközök beszerzése
2.4 Milyen gyors egy CUDA mag?
3 Ismerkedés a CUDA SDK-val
3.1 Hogyan lehet a CUDA-t programozni?
3.2 Mit csinál egy CUDA program?
3.3 CUDA projekt létrehozása
3.4 Dokumentáció
3.5 Kód vonatkozású alapfogalmak
3.6 A legegyszerűbb program
3.7 CUDA eszközök számának lekérdezése
3.8 Eszközadatok lekérdezése
4 Párhuzamos programozási alapok
4.1 Memóriaterületek kezelése, elérése
4.2 Tanácsok a CUDA kódoláshoz
4.3 Szálak és blokkok
4.4 Párhuzamosítási technikák
4.4.1 Navigálás a végrehajtási szálakon belül
4.5 Egyszerű FOR ciklus, CUDA nélkül
4.6 Egyszerű FOR ciklus, CUDA-val
4.7 Párhuzamosított FOR ciklus, CUDA-val
4.7.1 Több szál megadása
4.7.2 Több szál és blokk megadása
4.7.3 Szinkronizálás
4.8 CUDA programok debuggolása
4.8.1 Hibakeresés az NSight beépülő modullal
4.8.2 CUDA hibák kiírása
5 Haladó CUDA programozás
5.1.1 Több videokártya használata
5.1.2 CUDA programok teljesítménymérése
5.1.3 Futtatási profilok elemzése
6 Függelék
6.1 CUDA alapú renderlő algoritmusok
6.1.1 Pufferelés megvalósítása
6.1.2 Elemi alakzatok kirajzolása CUDA kernelekkel
6.2 A könyvben használt CUDA függvények
1 Előszó
Az NVIDIA napjaink egyik legmeghatározóbb technológiai cége, élen jár a mesterséges intelligencia kutatásában, a játékipart kiszolgáló grafikus technológiák kifejlesztésében és még számos más területen is. A cég egyik alappillére a CUDA technológia, mely a nagyközönség számára is elérhető, már a legolcsóbb NVIDIA videokártyákban is. Szinte fillérekért valóságos kis szuperszámítógépeket vihetünk haza.
A CUDA technológia a nagyteljesítményű párhuzamos adatfeldolgozásról szól, melyet számos területen fel lehet használni, például grafikus megjelenítéshez is.
Ez a leírás a CUDA platform programozásába nyújt bevezetést, egyszerű, közérthető formában.
1.1 Célközönség, követelmények
Az elsődleges célközönség a C programozásban jártas olvasó, aki rendelkezik tapasztalattal Windows alkalmazások fejlesztésében.
A Visual Studio mint fejlesztőeszköz alapszintű ismerete elengedhetetlen.
1.2 Letölthető programkódok
A leírásban található kódokat használó példaalkalmazások teljes forráskódja letölthető a következő weboldalról:
A leírás néhány grafikus példakódjának megalkotásakor ezen leírás szerzője sok inspirációt merített Dmitry V. Sokolov tinyrenderer leírásaiból: https://github.com/ssloy/tinyrenderer
2 A CUDA platform bemutatása
2.1 Mi az a CUDA?
Mielőtt fejest ugranánk a kódolásba, megismerkedünk azzal, hogy mi is egyáltalán a CUDA és hol helyezkedik el a technológia világában.
A CUDA a Compute Unified Device Architecture rövidítése, mely magyarul kb. annyit tesz: Egységesített Számítási Eszközarchitektúra.
A CUDA segítségével általános feladatokat gyorsíthatunk fel, CPU-alapú megközelítésekhez képest sokszoros mértékben.
A CUDA első megjelenése 2006-ra nyúlik vissza, a GeForce 8800-as kártyával kezdődően, mely a világon elsőként kínált általánosan programozható feldolgozóegységeket, közismertebb nevükön CUDA magokat.
Az NVIDIA cég fejlesztései olyannyira innovatívak és olcsók, hogy néhány év alatt az ipar egyik éllovasa lett és megkerülhetetlen szereplővé vált globálisan is. Jensen Huang, a cég egyik ikonikus alapítója számos rendezvényen megjelenik, bemutatva a legújabb fejlesztéseket.
Az NVIDIA a CUDA-t manapság platformként igyekszik pozicionálni, elsősorban a mesterséges intelligencia, a nagyteljesítményű számításokat igénylő alkalmazások, az autonóm eszközök és természetesen a grafikus megjelenítés köré szervezve.
Valóban, rengeteg alkalmazás, könyvtár és robbanásszerűen növekvő fejlesztői közösség jellemzi ezt az ökoszisztémát.
Az alábbi kép az NVIDIA által propagált bemutatása az ún. CUDA-X platformnak.
A CUDA architektúrák hagyományosan valamilyen híres tudósról kapják a nevüket.
Az egyes architektúrák elnevezése, időben:
• Tesla
• Fermi
• Kepler
• Maxwell
• Pascal
• Volta
• Turing
• Ada
• Hopper
• Blackwell.
Egy-egy új architektúrát jellemzően pár éves időközökben hoznak ki. Általában minden architektúra esetén rendre GeForce és Quadro kártyák is készülnek, bár ez nem törvényszerű.
2.2 CUDA alapú eszközök
Az ökoszisztémát az elérhető eszközök csoportjaival is szemléltethetjük.
Quadro RTX 8000: csúcskategóriás professzionális videokártya.
GPU órajele: 1440 MHz
CUDA® magok száma: 4608
Memória órajele: 14000 MHz
Memória mérete: 48 GB
Memória busz: 384 bit
Memória sávszélessége: 672 GB/s
Quadro P420: belépő szintű, de professzionális felhasználásra kifejlesztett videokártya (Pascal architektúra).
GPU órajele: 1228 MHz
CUDA® magok száma: 256
Memória órajele: 4008 MHz
Memória mérete: 2 GB
Memória busz: 64 bit
Memória sávszélessége: 32 GB/s
Geforce RTX 2080 Ti: csúcskategóriás játékkártya.
GPU órajele: 1350 MHz
CUDA® magok száma: 4352
Memória órajele: 14000 MHz
Memória mérete: 11 GB
Memória busz: 352 bit
Memória sávszélessége: 616 GB/s
GT1650 SUPER: alsó-középkategóriás játékkártya
GPU órajele: 1530 MHz
CUDA® magok száma: 1280
Memória órajele: 12000 MHz
Memória mérete: 4 GB
Memória busz: 128 bit
Memória sávszélessége: 192 GB/s
Tesla T4: kifejezetten gépi tanulásra kifejlesztett gyorsítókártya.
GPU órajele: 585 MHz
CUDA® magok száma: 2560
Memória órajele: 5001 MHz
Memória mérete: 16 GB
Memória busz: 256 bit
Memória sávszélessége: 320 GB/s
SHIELD TV PRO: otthoni szórakoztató elektronikai eszköz és videojáték konzol.
CUDA® magok száma: 256
Memória mérete: 3 GB
JETSON AGX Xavier: autonóm gépek irányításához tervezett processzor.
CUDA® magok száma: 512
Memória mérete: 32 GB
DRIVE AGX Orin és Pegasus: önvezető járművek vezérlésére kifejlesztett, rendkívül komplex rendszerek.
2.3 CUDA eszközök beszerzése
CUDA eszközökhöz legegyszerűbben számítástechnikai kereskedésekben juthatunk videokártyák formájában.
A kártyák beszerzésekor a következő dolgokra kell figyelni:
• Egy, vagy több kártya szimultán programozását tervezzük-e: a számítógép elérhető alaplapi csatlakozónyílásait ellenőrizzük!
• Mennyit fogyaszt a kártya: adott esetben egy erősebb tápegység beszerzése is szükséges lehet.
• Fontos-e a kiugróan nagy teljesítmény: ha nem, akkor tanuláshoz, ismerkedéshez a legolcsóbb videokártya is megteszi, mivel ugyanaz a fejlesztési elv használható velük is, mint a legdrágább kártyákkal. A későbbiekben programjaink egy esetlegesen beszerzett nagyobb teljesítményű kártyán is ugyanúgy fognak tudni futni, bár megjegyzendő, hogy a legolcsóbb kártyák is érezhetően nagy teljesítményt képesek nyújtani.
Ne feledjük, hogy egy alkalmazás adatfeldolgozási sebessége hardveres és szoftveres feltételektől is függ.
A szoftveres feltétel az algoritmus és a vonatkozó keretrendszer, amelyeknek minél hatékonyabbnak kell lenniük és minél jobban ki kell tudniuk használni a platformot.
Az alábbi hardveres feltételeket is érdemes figylemebe vennünk:
• CPU teljesítménye
• RAM sebessége
• alaplapi csatolók, buszok sebessége
• CUDA eszköz teljesítménye
• adathálózat átviteli sebessége (online rendszereknél)
• háttértároló sebessége.
Ideális esetben minden hardvereszköz jó minőségű és optimális összhangban van, tehát nem kell indokolatlanul sokat várniuk egymásra.
A leírás példáinak megírásához egy Intel i5 CPU-t, 16GB RAM-ot, SSD meghajtót, 2 darab Geforce GTX 1650 SUPER videokártyát és egy 650W teljesítményű tápegység állt rendelkezésre.
2.4 Milyen gyors egy CUDA mag?
Egyetlen CUDA mag sebessége nem felel meg egy CPU magénak (megjegyezzük, hogy a többmagos CPU-k magjai sem ugyanazt a sebességet tudják produkálni).
A CUDA egy eleve párhuzamos feldolgozásra tervezett architektúra, amit saját magához, vagy legalábbis hasonló architektúrákhoz van értelme csak hasonlítani.
Ennek ellenére elmondható, egy GeForce 1030-as GPU teljesítménye nagyjából megfelel kb. két és fél darab Intel i5-ös processzor nyers teljesítményének (egy i5-ös processzor teljesítménye durván 0,4 TFLOP).
3 Ismerkedés a CUDA SDK-val
Ez a leírás a PC-k PCI Express bővítőhelyeibe illeszthető CUDA kártyák programozásához nyújt segítséget. Az ettől eltérő CUDA eszközökre történő fejlesztéshez külön SDK-kat kell beszereznünk.
3.1 Hogyan lehet a CUDA-t programozni?
A CUDA SDK innen tölthető le:
https://developer.nvidia.com/cuda-downloads
A CUDA SDK főbb komponensei:
• SDK a fordítóprogrammal
• dokumentáció
• példaprogramok
• Visual Studio integráció
• kiegészítő programozási könyvtárak, melyek szintén a CUDA-n alapulnak
• eszközmeghajtó.
A CUDA SDK-val feltelepíthető meghajtóprogram sokszor nem megfelelően működik együtt a CUDA SDK-val, érdemes mindig járulékosan telepíteni az NVIDIA legfrissebb vonatkozó meghajtóprogramját is.
CUDA esetében fontos fogalmak a CUDA SDK verziója és a CUDA Compute Capability verziója. Előbbi a fejlesztőkészletre vonatkozik, utóbbi pedig a fizikai eszköz számítási architektúráját jelenti, mely hardverfüggő érték.
A már a Turing architektúrát is támogató 10-es verziójú CUDA SDK példaprogramjait az alábbi könyvtárban fogjuk megtalálni:
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v10.0\
A CUDA alapértelmezett programozási nyelve a CUDA C, mely a sztenderd C nyelv kiterjesztéseként képzelhető el. A CUDA kódokat az NVIDIA saját fordítója fordítja le, de az egész folyamat integrálva van egy már meglévő fejlesztőrendszerbe, tipikusan a Visual Studio-ba.
3.2 Mit csinál egy CUDA program?
Egy nagyon fontos kérdést tisztáznunk kell: hogyan működik egy CUDA program?
A CUDA magok adatfeldolgozásra használhatóak: szorzás, osztás, stb., viszont mindezt extrém módon párhuzamosított formában lehet elérni.
A feldolgozandó adatokat minden esetben a CPU alapú kódrészeknek kell előkészíteniük, majd az összes releváns adatot át kell másolni a CUDA kártya memóriájába és ott csak ezután kezdődhet el a feldolgozás. Az esetleges eredményeket, adatokat valamilyen formában szintén vissza kell másolnunk a számítógép fő memóriájába, hogy a programunk valamit kezdhessen ezzel, például kiírja a képernyőre.
A főbb fázisok tehát:
• adatok előpkészítése
• adatmásolás a CUDA eszközre
• adatfeldolgozás
• adatvisszamásolás (eredmények).
Programozási szempontból kritikus az adatmásolások sebessége.
3.3 CUDA projekt létrehozása
Feltelepített CUDA SDK-val a Visual Studio projektvarázslója automatikusan tartalmazni fog egy NVIDIA/CUDA bejegyzést, a konkrét CUDA SDK verziószámával.
Egyszerűen csak ki kell választanunk és így létrehoznunk egy új projektet!
A projektsablon tartalmaz egy egyszerű CUDA-alapú konzolalkalmazást, amit természetesen módosíthatunk.
Nem kell megijedni, ha a Visual Studio ún. IntelliSense funkciója bizonyos kódrészeket pirossal aláhúz! CUDA projektek esetében ez akkor is előfordul, ha a programkód egyébként tökéletes szintaktikailag. Ez normális.
3.4 Dokumentáció
A CUDA programozás egy angol nyelven elérhető kiváló referenciagyűjteménnyel rendelkezik.
Központi oldal, SDK letöltése:
https://developer.nvidia.com/cuda-zone
A teljes, hivatalos dokumentáció elérhetősége:
https://docs.nvidia.com/cuda/index.html
Támogatott eszközök:
https://developer.nvidia.com/cuda-gpus
Fejlesztőkészlet, dokumentáció:
https://developer.nvidia.com/cuda-downloads
Bizonyos eszközökhöz, dokumentációkhoz egy ingyenesen beregisztrálható NVIDIA fejlesztői fiókra lehet szükségünk.
3.5 Kód vonatkozású alapfogalmak
A legfontosabb alapfogalmak a következők.
Host (CPU) kód: a program azon része, melyet a CPU hajt végre.
Device (eszköz/GPU) kód: a programunknak azon része, melyet a GPU, a CUDA magokkal hajt végre.
Kernel: azok a konkrét függvények (device kódok, függvényhívások), melyeket a CUDA C fordítója fordít és a CUDA eszköz hajt végre.
Végrehajtási konfiguráció: annak definiálása, hogy egy kernel hány ún. végrehajtási blokkon és szálon fusson le. A megadás <<< és >>> jelek közé adandó meg, számok formájában.
A megfelelő konfiguráció kritikus a teljesítmény szempontjából. A minimálisan megadandó érték <<< 1, 1 >>>.
Akárcsak CPU-k esetében, úgy a GPU-k esetében is lehetséges több szálon futtatni kódokat, de CUDA esetében ez akár több millió logikai végrehajtási szálat is jelenthet.
Egy időben annyi végrehajtási szál tud futni, amennyi CUDA maggal rendelkezik az adott eszköz.
3.6 A legegyszerűbb program
Minden CUDA program .CU kiterjesztéssel rendelkezik. A Visual Studioban a CUDA projektsablonnal létrehozott főprogramunk forrásfájljának neve alapértelmezetten KERNEL.CU lesz.
A legegyszerűbb CUDA program ugyanúgy néz ki, mint egy közönséges C nyelvű alkalmazás (nyugodtan írjuk felül az alapértelmezett kódsablont).
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
int main()
{
printf("Elso CUDA programom...");
return 0;
}
Sikeres fordítás után egy .EXE fájl fog létrejönni, melyet természetesen önállóan is futtathatunk. Mindemellett, ez a program természetesen még semmilyen GPU-specifikus dolgot nem csinál.
Mielőtt kifejezetten a párhuzamosítás gyakrolatát kezdenénk elsajátítani, néhány alapvető CUDA-lekérdezés elvégzését is meg kell ismernünk.
3.7 CUDA eszközök számának lekérdezése
Nem árt tudni, hogy hány CUDA eszköz található egy számítógépben, arra az esetre, ha esetleg el szeretnénk osztani számításokat közöttük.
Egy ilyen lekérdezésre mutat példát az alábbi rövidke kis program (CUDA_01.CU):
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
int main()
{
int darabszam;
cudaGetDeviceCount(&darabszam);
printf("%i darab CUDA eszkozt talaltam...",darabszam);
return 0;
}
A program kimenete:
2 darab CUDA eszkozt talaltam...
3.8 Eszközadatok lekérdezése
További hasznos gyakorlati módszer a CUDA eszközök legfontosabb technikai tulajdonságait lekérdezni. Például, hogy milyen verziójú számítási technológiát támogat eszközünk, vagy hány CUDA mag található az eszközön.
Az alábbi példa lefordításához meg kell adnunk a helper_cuda.h fejlécállomány pontos helyét is. A Visual Studio-ban a projekt tulajdonságai párbeszédablakban a VC++ Directories alatt tudjuk megadni az elérési utat.
C:\ProgramData\NVIDIA Corporation\CUDA Samples\v10.0\common\inc;
Nem árt észben tartani, hogy a CUDA fejlécállományok másik pontos „lelőhelye” az alábbi könyvtár:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.0\include
Egy komplex példakód (CUDA_02.CU) eszköztulajdonságok lekérdezéséhez:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include "helper_cuda.h"
#include <stdio.h>
int main()
{
int darabszam;
int dev, driverVersion = 0, runtimeVersion = 0;
cudaGetDeviceCount(&darabszam);
if (darabszam == 0) printf("Nem talaltam tamogatott CUDA eszkozt!\n");
else printf("%i darab CUDA eszkozt talaltam...", darabszam);
for (dev = 0; dev < darabszam; ++dev)
{
cudaSetDevice(dev);
cudaDeviceProp deviceProp;
cudaGetDeviceProperties(&deviceProp, dev);
printf("\n%d. eszkoz neve: \"%s\"\n", dev, deviceProp.name);
cudaDriverGetVersion(&driverVersion);
cudaRuntimeGetVersion(&runtimeVersion);
printf(" CUDA meghajto verzioja / Futtato kornyezet verzioja %d.%d / %d.%d\n",
driverVersion / 1000, (driverVersion % 100) / 10,
runtimeVersion / 1000, (runtimeVersion % 100) / 10);
printf(" CUDA Capability verzio: %d.%d\n",
deviceProp.major, deviceProp.minor);
char msg[256];
sprintf_s(msg, sizeof(msg),
" Osszmemoria: %.0f MBytes "
"(%llu bytes)\n",
static_cast<float>(deviceProp.totalGlobalMem / 1048576.0f),
(unsigned long long)deviceProp.totalGlobalMem);
printf("%s", msg);
printf(" (%2d) Multiprocesszor, (%3d) CUDA mag/MP: %d CUDA mag\n",
deviceProp.multiProcessorCount,
_ConvertSMVer2Cores(deviceProp.major, deviceProp.minor),
_ConvertSMVer2Cores(deviceProp.major, deviceProp.minor) *
deviceProp.multiProcessorCount);
printf(
" GPU max. orajel: %.0f MHz (%0.2f "
"GHz)\n",
deviceProp.clockRate * 1e-3f, deviceProp.clockRate * 1e-6f);
printf(" Memoria orajel: %.0f Mhz\n",
deviceProp.memoryClockRate * 1e-3f);
printf(" Memoria atviteli sebesseg: %d-bit\n",
deviceProp.memoryBusWidth);
printf(" Warp merete: %d\n",
deviceProp.warpSize);
printf(" Szalak max. szama / multiprocessor: %d\n",
deviceProp.maxThreadsPerMultiProcessor);
printf(" Szalak max. szama / blokk: %d\n",
deviceProp.maxThreadsPerBlock);
printf(" Blokk maximalis dimenzioi (x,y,z): (%d, %d, %d)\n",
deviceProp.maxThreadsDim[0], deviceProp.maxThreadsDim[1],
deviceProp.maxThreadsDim[2]);
printf(" Grid maximalis dimenzioi (x,y,z): (%d, %d, %d)\n",
deviceProp.maxGridSize[0], deviceProp.maxGridSize[1],
deviceProp.maxGridSize[2]);
}
return 0;
}
A program kimenete a következő lesz:
2 darab CUDA eszkozt talaltam...
0. eszkoz neve: "GeForce GTX 1650 SUPER"
CUDA meghajto verzioja / Futtato kornyezet verzioja 10.2 / 10.0
CUDA Capability verzio: 7.5
Osszmemoria: 4096 MBytes (4294967296 bytes)
(20) Multiprocesszor, ( 64) CUDA mag/MP: 1280 CUDA mag
GPU max. orajel: 1725 MHz (1.73 GHz)
Memoria orajel: 6001 Mhz
Memoria atviteli sebesseg: 128-bit
Warp merete: 32
Szalak max. szama / multiprocessor: 1024
Szalak max. szama / blokk: 1024
Blokk maximalis dimenzioi (x,y,z): (1024, 1024, 64)
Grid maximalis dimenzioi (x,y,z): (2147483647, 65535, 65535)
1. eszkoz neve: "GeForce GTX 1650 SUPER"
CUDA meghajto verzioja / Futtato kornyezet verzioja 10.2 / 10.0
CUDA Capability verzio: 7.5
Osszmemoria: 4096 MBytes (4294967296 bytes)
(20) Multiprocesszor, ( 64) CUDA mag/MP: 1280 CUDA mag
GPU max. orajel: 1725 MHz (1.73 GHz)
Memoria orajel: 6001 Mhz
Memoria atviteli sebesseg: 128-bit
Warp merete: 32
Szalak max. szama / multiprocessor: 1024
Szalak max. szama / blokk: 1024
Blokk maximalis dimenzioi (x,y,z): (1024, 1024, 64)
Grid maximalis dimenzioi (x,y,z): (2147483647, 65535, 65535)
4 Párhuzamos programozási alapok
Ebben a fejezetben megismerjük a CUDA alapú párhuzamos programozás alapjait. Ennek a tudásnak a birtokában saját CUDA kernelek megírására lesz képes az olvasó.
Megjegyezzük, hogy a CUDA programozás igen szerteágazó, számos haladó témát nem tárgyalunk, mint például a CUDA stream-ek írását, különböző GPU memóriatárolók használatát, atomi műveleteket, stb.
Ne aggódjunk azonban, mert ezen technikák ismerete nélkül is messzemenően ki fogjuk tudni használni a CUDA eszközünk lehetőségeit!
4.1 Memóriaterületek kezelése, elérése
Fontos tisztázni, hogy a videokártyánk hogyan fér hozzá az adatokhoz. A mai grafikus gyorsítók memóriájának sebessége túlszárnyalja az asztali gépekben található RAM sebességét. A kereskedelmi forgalomban elérhető kártyák között jelenleg a 14000MHz sebességű memóriával rendelkezők a leggyorsabbak. Ezzel szemben az asztali gépekhez 4000MHz-nél gyorsabb RAM nem szokott előfordulni. A memória sebessége nem minden, de igen fontos faktor, ahogyan arra a 2.3 fejezetben is utaltunk.
Miután programunk már tudja, hogy milyen adatokat akar feldolgozni CUDA-val, le kell foglalnunk a szükséges memóriaterületet az eszközön. A CUDA C tartalmaz egy cudaMalloc függvényt, ami a C nyelv klasszikus malloc memóriafoglaló függvényéhez hasonlóan használható.
A videokártya teljes szabad memóriakapacitását kihasználhatjuk a foglaláskor. A gyakorlatban ez valamennyivel kevesebb mint a teljes memória, mivel az operációs rendszer mindenképpen lefoglal 2-300 MB területet.
A memóriaterületet egy mutatóval foglalhatjuk le és ezen keresztül is érhetjük el.
Példa memória foglalásra:
unsigned int *mutato;
cudaMalloc((void**)&mutato, 1000000 * sizeof(unsigned int));
Ezután a tényleges adatokat be kell másolnunk az asztali gép RAM-jából a CUDA eszköz memóriájába. Ehhez a cudaMemcpy függvényt használhatjuk, melynek utolsó paramétere a memóriamozgatás irányát szabja meg.
A legfontosabb irányok:
• cudaMemcpyHostToDevice: Host -> Device
• cudaMemcpyDeviceToHost: Device -> Host
• cudaMemcpyDeviceToDevice: Device -> Device.
Fontos, hogy a forrás- és a célmemóriaterület ugyanolyan típusú adatokat tartalmazzon és hogy a célmemóriaterület mérete legalább akkora legyen, mint a forrásé.
Az adatmásolás után a CUDA kernelek műveleteket végezhetnek az adatokon.
A CUDA programok futásának befejezésekor a lefoglalt memóriaterületeket az eszközön is fel kell szabadítanunk. Erre szolgál a cudaFree függvény, mely a sztenderd C nyelv free függvényéhez hasonlóan működik.
A memóriafelszabadítás ablakos alkalmazás esetében a WM_CLOSE eseménykezelőben lesz ideális:
case WM_CLOSE:
pRT->Release();
pD2DFactory->Release();
cudaFree(mutato);
DestroyWindow(hwnd);
return 0;
4.2 Tanácsok a CUDA kódoláshoz
Az adatok másolgatása nem öncélú dolog, hanem szükséges kényszerűség is, ugyanis a CUDA eszközök egyszerűen nem férhetnek hozzá a RAM-hoz és viszont. Ezt olyannyira szigorúan kezeli az SDK, hogy a kód le sem fordítható, ha ilyen nem megengedett memóriahivatkozásokkal próbálkoznánk. A helytelen memóriahozzáférés különben az egyik leggyakoribb hibaforrás.
További fontos tudnivaló, hogy a kernelek futása közben fellépő hibák a kernel futásának azonnali megszakítását szokták maguk után vonni. Sokszor segít, ha a felhasznált CUDA eszközt a cudaReset függvénnyel alaphelyzetbe állítjuk. Ezután viszont gondoskodnunk kell az összes objektum, például memóriaterületek stb. ismételt létrehozásáról.
Kernelek indításakor ajánlott a függvények paramterlistájában átadni a feldolgozandó adatok mennyiségét na és persze a rájuk mutató mutatót is. Így a kernelen belül végig kézben tarthatjuk az adatfeldolgozás folyamatát.
Példa kernel paraméterezésére:
__device__ void pelda_kernel(int maxelemszam, unsigned int *adatok)
CUDA eszközön használt változókat memóriafoglalás nélkül is deklarálhatunk főprogramunkban, ehhez a ___device__ kulcsszót kell a deklarációban elé írnunk.
Érdekesség, hogy amennyiben nem sok adatot, hanem csak egy-egy értéket adunk át a kernelnek, akkor a változó lehet a RAM-ban is, vagy a videokártya memóriájában: az értékét mindenképpen felhasználhatjuk a kernelben. Fordítva viszont ez nem megengedett: a CUDA eszközön létrehozott változóra nem hivatkozhatunk host kódból!
Az a legbiztosabb, ha minden memóriaterületet dinamikusan foglalunk le.
4.3 Szálak és blokkok
CUDA esetében a GPU ún. SIMT (Single Instruction, Multiple Thread) elven működik:. minden végrehajtási szál ugyanazt a kódrészt/utasítást hajtja végre.
Egy kernel párhuzamosított végrehajtása blokkokba rendezett végrehajtási szálakon megy végbe. A blokkok és szálak száma alkotja az ún. végrehajtási konfigurációt.
Egy blokk adott számú szálat foglalhat magában maximálisan. A blokkok összessége a grid.
Egy blokkon belül a konkrét végrehajtási szálakat az egyes CUDA magok dolgozzák fel.
A végrehajtási konfigurációt speciális módon kell megadnunk, <<< A , B >>> jelek között, számokkal. Az A érték a blokkok számát jelenti, a B a szálakét. Lássunk néhány példát!
Az <<< 1 , 1 >>> végrehajtási konfiguráció „klasszikus” egyszálas kódvégrehajtást eredményez.
Az <<< 1, 100 >>> egy blokkot ad meg, 100 végrehajtási szállal.
Az <<< 100, 1 >>> 100 blokkot ad meg, melyek mindegyike 1 végrehajtási szálat fog lefuttatni.
Mint látható, rengeteg féle módon beállíthatjuk a konfigurációt. A dolog lényege éppen a finomhangolhatóságban van. Közel sem biztos ugyanis, hogy mindig az összes magon / szálon akarunk kerneleket futtatni. Jó példa erre a procedurális képgenerálás, ahol nem „agyatlanul” a maximális blokk- és szál megadása a cél. Pontosan ezért is jó, hogy le tudjuk kérdezni egy adott eszközön használható maximális értékeket a végrehajtási konfigurációhoz
A CUDA architektúra az ún. Streaming Multiprocessor-ok (SMP) köré van felépítve. Egy SMP meghatározott számú CUDA magot tartalmaz. A konkrét párhuzamos végrehajtásokat az SMP-k vezénylik. A CUDA magok száma SMP architektúránként, vagy ritkábban GPU-ként is változik. A Pascal esetében például egy SMP-hez 128 CUDA mag tartozik. Turing architektúra esetében ez 64.
Meg kell említeni, hogy az NVIDIA GPU-k a háttérben ún. szálcsoportokat hajtanak végre. Egy szálcsoport 32 szálból áll és a neve WARP. Nyitva áll annak a lehetősége, hogy a CUDA programozó WARP-szinten befolyásolja a párhuzamos végrehajtásokat.
4.4 Párhuzamosítási alaptechnikák
Hogyan képzelhető el józan paraszti ésszel egy párhuzamos végrehajtás? A legegyszerűbb úgy elképzelni a dolgot, mintha mindegyik végrehajtási szál az adott kernel egy külön másolatát futtatná.
A végrehajtás olyannyira párhuzamos, hogy nem is garantált az, hogy a szálak milyen sorrendben futnak le! A nagy teljesítmény egyik ára, hogy a végrehajtást a hardverre bízzuk.
Ha több blokkot és szálat adunk meg, mint amit egyszerre kínál az eszköz, akkor a hardver újra és újra kiosztja a feladatokat a CUDA magokra, amíg marad futtatandó szál.
Fontos, hogy minden szál csak egyszer fut le, nincsen ismétlődés! Ha egy szál lefut, akkor nem indítható el ismét.
A fentiekből következik, hogy a kerneleket önálló futási egységekként kell megírni.
Nézzünk egy példát! Ha 2 blokkot és blokkonkét három szálat adunk meg, akkor összesen 6 különböző szál fog lefutni egyidőben, a hardver által meghatározott sorrendben:
1. blokk – 1. szál
1. blokk – 2. szál
1. blokk – 3. szál
2. blokk – 1. szál
2. blokk – 2. szál
2. blokk – 3. szál
A hardver tehát a végrehajtás sorrendjét nem, de a megtörténtét garantálni tudja. Ezt egy szinkronizációs függvényhívással (lásd 4.7.3. fejezetet) várhatjuk meg, mert meg kell tudnunk mondani, hogy minden szál befejezte-e a futását.
A klasszikus programozási nyelvek sok adatot ciklusokkal, például FOR végrehajtási ciklussal dolgoznak fel.
CUDA-ban sok adatot FOR ciklusokra támaszkodva is feldolgozhatunk, de akár teljesen ki is válthatjuk ezt. Előző esetben a végrehajtási konfigurációból levezetett léptéket állítunk be a ciklusban.
Ez a leírás preferáltan a ciklusos megközelítést mutatja be a példákban.
4.4.1 Navigálás a végrehajtási szálakon belül
Ismernünk kell, hogy miképpen lehet ellenőrizni a kernel végrehajtásakor, hogy éppen melyik szál hajtja végre azt és melyik blokkban.
A CUDA ehhez globális változókat használ, egyszerűen csak le kell kérdeznünk őket.
Ezek a következők:
blockIdx.x: az aktuális blokk számát adja vissza.
gridDim.x: az összes blokk számát adja vissza.
ThreadIdx.x: az aktuális szál számát adja vissza.
blockDim.x: szálak összes száma egy blokkban.
Miért szerepel a változók utáni .x tulajdonság? Azért, mert mind a blokkokat, mind a szálakat több dimenzióban is megadhatjuk, akár három dimenzióban is. Amikor ezt adjuk meg <<< 1 , 1 >>>, akkor valójában ez lesz beállítva:
<<< (1 , 1, 1) , ( 1, 1, 1 ) >>>
A több dimenzió megadása így történhet explicit módon: dim3(1,1,1).
Ez a fajta felosztás lehetőséget ad extrém nagy mennyiségű adat párhuzamos feldolgozásának a megszervezésére.
4.5 Egyszerű FOR ciklus, CUDA nélkül
Az alábbiakban egy hagyományos, egy szálon futó példa (CUDA_03.CU) látható, a CPU_kernel függvény köré kialakítva. Ezt a függvényt fogjuk a későbbiekben átírni CUDA kernellé.
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
void CPU_kernel(int maxelemszam)
{
int i;
for (i = 0; i < maxelemszam; ++i)
{
printf("%i\n", i);
}
}
int main(void)
{
CPU_kernel(100);
printf("Vegrehajtas befejezve!\n");
return 0;
}
A program kimenete:
89
90
91
92
93
94
95
96
97
98
99
Vegrehajtas befejezve!
4.6 Egyszerű FOR ciklus, CUDA-val
Végre lássunk egy valódi CUDA kernelt tartalmazó példát (CUDA_04.CU) is! Mindazonáltal itt sem a kernel, sem annak végrehajtása még semmiféle párhuzamosítást nem tartalmaz!
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void GPU_kernel(int maxelemszam)
{
int i;
for (i = 0; i < maxelemszam; ++i)
{
printf("%i\n", i);
}
}
int main(void)
{
GPU_kernel << < 1, 1 >> > (100);
cudaDeviceSynchronize();
printf("Vegrehajtas befejezve!\n");
return 0;
}
A program kimenete nagyon hasonló a CPU-s változatéhoz:
89
90
91
92
93
94
95
96
97
98
99
Vegrehajtas befejezve!
A példaprogramunkban egyetlen egy blokk és szál futott le, ami lényegében megfelel az egyszálas CPU kód végrehajtásának.
A programkódban számos különbséget felfedezhetünk. Az első, egyik legfontosabb különbség, hogy a CUDA kernelek elé oda kell írnunk a __global__, vagy a __device__ kulcsszavakat. Előbbi azt jelzi, hogy a függvény a GPU-n fog futni és csak a CPU kódból hívható meg. A __device__ kulcsszó pedig azt jelzi, ha a kernel-t csak egy másik CUDA kernelből lehet meghívni.
A fentieket könnyen leellenőrizhetjük például úgy is, hogy egy __device__ függvényt CPU kódból próbálunk meghívni. Fordítási hibát fogunk kapni a fordítás során.
További fontos különbség a GPU kernel indításának kezdeményezése is. A meghívásakor a függvény neve után <<< >>> jelek között kell megadnunk a végrehajtási konfigurációt, csak utána szerepeltethetjük a függvény paraméterlistáját, ha van.
4.7 Párhuzamosított FOR ciklus, CUDA-val
Most megnézzük, miként írhatjuk át a GPU kódot valódi párhuzamos végrehajtásra alkalmas CUDA kernellé.
4.7.1 Több szál megadása
Először 1 blokk és 10 szál kezelését nézzük meg egy példán (CUDA_05.CU) keresztül, ahol a végrehajtási konfiguráció értelemszerűen ez: <<< 1 , 10 >>>.
A kernelt a következőképpen kell megadnunk a forráskódban, a program forráskódja egyébként mindenben ugyanaz marad, mint az előzőekben:
__global__ void GPU_kernel(int maxelemszam)
{
int i;
int startindex = threadIdx.x;
int leptek = blockDim.x;
for (i = startindex; i < maxelemszam; i += leptek)
{
printf("%i\n", i);
}
}
Kritikus annak ellenőrzése, hogy az adott szál sorszáma nem lépi-e túl a maximális elemszámot.
A fenti függvény FOR ciklus nélküli változata egyébként így néz ki (CUDA_05_FOR_NELKUL.CU), függvényünk így még kompaktabb lehet:
__global__ void GPU_kernel(int maxelemszam)
{
int aktualis_index = threadIdx.x + (blockIdx.x * blockDim.x);
if (aktualis_index < maxelemszam) printf("%i\n", aktualis_index);
}
4.7.2 Több szál és blokk megadása
Amennyiben több blokkot is használni kívánunk, ajánlott azt a maximális blokk-elemszámból és a szálak tervezett számából kiszámolni. Ezt egy ún. grid-stride cikluson keresztül mutatjuk be.
Ha 128 szálat akarunk indítani 100.000 elem feldolgozásához, akkor a blokkok ideális számát az alábbi képlettel számolhatjuk ki (a gyakorlatban ezt azért érdemes kipróbálással tovább finomítani):
(MAX_ELEMSZAM + SZALAK_SZAMA - 1) / SZALAK_SZAMA
Programkóddal:
int szalak_szama = 128;
int blokkok_szama = (100000 + szalak_szama - 1) / szalak_szama;
A CUDA kernel kódja:
__global__ void GPU_kernel(int maxelemszam)
{
int i;
int startindex = threadIdx.x + (blockIdx.x * blockDim.x);
int leptek = blockDim.x * gridDim.x;
for (i = startindex; i < maxelemszam; i += leptek)
{
printf("%i\n", i);
}
}
A kernel elindítása így történhet:
GPU_kernel <<< blokkok_szama, szalak_szama >>> (100000);
A program (CUDA_06.CU) kimenete tanulságos, ugyanis a legutolsó kiírt érték nem a legnagyobb (999999) sorszám. Ez is jól példázza, mennyire önállóan futtatja a hardver a kerneleket:
999997
999998
999999
999424
999544
999545
999546
999547
999548
999549
Vegrehajtas befejezve!
4.7.3 Szinkronizálás
A kernelek végrehajtása aszinkron módon megy végbe, emiatt a programunk a CPU szálon futva nem várná meg, amíg az összes szál lefut, hanem automatikusan továbbhaladna. Ezt elkerülendő, a host kódban minden kernel indítás után explicit meg kell várni, amíg befejeződik a kernelek végrehajtása.
Erre szolgál a cudaDeviceSynchronize() függvényhívás, amit minden kernelhívás után meg kell adni!
Nem kell viszont kiadnunk CUDA memóriamásoló függvények meghívása után (például: cudaMemcpy), mivel azok szinkron módon hajtódnak végre.
Egy egyszerű példakód (CUDA07.CU):
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void GPU_kernel()
{
printf("blokk: %i, szal:%i\n", blockIdx.x, threadIdx.x);
}
int main(void)
{
GPU_kernel << < 10, 2 >> > ();
cudaDeviceSynchronize();
printf("Vegrehajtas befejezve!\n");
return 0;
}
A program kimenete:
blokk: 6, szal:0
blokk: 6, szal:1
blokk: 0, szal:0
blokk: 0, szal:1
blokk: 9, szal:0
blokk: 9, szal:1
blokk: 4, szal:0
blokk: 4, szal:1
blokk: 3, szal:0
blokk: 3, szal:1
blokk: 7, szal:0
blokk: 7, szal:1
blokk: 1, szal:0
blokk: 1, szal:1
blokk: 5, szal:0
blokk: 5, szal:1
blokk: 8, szal:0
blokk: 8, szal:1
blokk: 2, szal:0
blokk: 2, szal:1
Vegrehajtas befejezve!
4.8 CUDA programok debuggolása
Mivel a CUDA programrészek a GPU-n futnak, nagyon fontos megismerkednünk azzal, hogy miképpen végezhetünk hibakeresést ezeken a kódokon, hiszen a sem a CPU, sem a GPU nem „látja” közvetlenül a másik fél memóriájának tartalmát.
4.8.1 Hibakeresés az NSight beépülő modullal
Az NSight debugger segédprogramot a CUDA SDK-val telepíthetjük és beépülő modulként használható.
A Visual Studio főmenüjében az 'NSight' menüpont alatt érhető el. A hibakereső futtatást a StartNSight Debugger menüpontot kiválasztva kezdeményezhetjük.
Sajnos azt nem lehet befolyásolni, hogy éppen melyik szálon történő végrehajtásnál fog megállni a kód futása debuggolás céljából.
4.8.2 CUDA hibák kiírása
CUDA programok készítése és tesztelése során a leggyakoribb hibák közé tartoznak a helytelen memóriacím hivatkozások.
Hibák felismeréséhez és kezeléséhez használható a CudaGetLastError() függvény, mely a legutolsó hiba rövid leírását szövegesen is vissza tudja adni.
A legutolsó hiba szövegét az alábbi módon kérdezhetjük le:
char hibauzenet[768];
strcpy_s(hibauzenet, cudaGetErrorString(cudaGetLastError()));
A CUDA függvények általában visszatérési értékükön keresztül is jelzik a végrehajtás sikerességét, vagy sikertelenségét.
Erre egy példa:
cudaError_t hiba = cudaMalloc(&adat, elemszam * sizeof(int));
if (hiba != cudaSuccess) printf("%s\n", cudaGetErrorString(err));
Ehhez hasonló logikát követ a kernelek végrehajtása során fellépő hibák lekérdezése is:
cudaError_t errAsync = cudaDeviceSynchronize();
if (errAsync != cudaSuccess) printf("Async kernel error: %s\n", cudaGetErrorString(errAsync));
5 Haladó CUDA programozás
Ebben a fejezetben olyan technikákat lehet találni, melyek az előző fejezetek ismeretanyagára építenek, de segítségükkel hatékonyabbá tehetjük programjainkat, illetve a fejlesztési munkát.
5.1.1 Több videokártya használata
Amennyiben több CUDA videokártya is van a gépünkben, jogos az igény arra, hogy elosszuk a munkát közöttük, vagy éppen kétszer-háromszor annyi adatot dolgoztassunk fel ugyanannyi idő alatt. Mindkét megközelítés lehetséges.
Az egyik legegyszerűbb módszer, ha a host kódból elindítjuk az egyik videokártyán a munkát, majd a cudaSetDevice() függvénnyel átváltunk a másik kártyára és ott is elindítjuk a számítási feladatokat. A cudaDeviceSynchronize( függvény pedig bevárja, amíg mindkét kártya befejezi a munkát.
Több videokártya használata drámaian ki tudja tolni lehetőségeinket, ráadásul bármilyen módszert és variációt megvalósíthatunk, csak a fantáziánk szabhat határt. Elképzelhető például olyan képgenerálás, mely a képkockákat elosztva rendereli a videokártyákon: a kép egyik felét az egyiken, a másikat a másikon. Ilyen esetekben gondoskodni kell arról, hogy az egyes CUDA eszközök eredményeit valamilyen módon összefésüljük vagy az egyik eszközön, vagy a CPU-val, a host oldalon.
Több eszköz esetén akkor a legkönyebb a dolgunk, ha mindegyik teljesen ugyanaz a modell, magyarul az összes kártya ugyanolyan. A dolog akkor lesz kacifántos, ha eltérő teljesítményű modellekkel kell dolgozni. Ilyenkor ugyanis gyorsulást csak akkor érhetünk el a feladatok elosztásával, ha a kisebb teljesítményű kártyára arányosan kevesebb feladatot osztunk ki. Máskülönben egy ilyen konfiguráció mindig a leglassabb kártyára fog várni és csak annyira lesz gyors, amennyire a leglassabb kártya, hiszen a programnak be kell várnia az összes kártyán levő munka befejeztét.
Ezt leszámítva azonban valóban arányos gyorsulásokat lehet elérni ezzel a módszerrel.
Az alábbi példaprogram (CUDA_08.CU) két CUDA videokártyával használható. Mindkét kártyán létrehozunk egy-egy int típusú változót, majd beállítjuk az egyik értékét 1-re, a másikét 9-re. Ezután egymás után átmásoljuk a videokártyákon tárolt értékeket az akt_adat nevű változóba és kiírjuk a képernyőre. A videokártyák között a cudaSetDevice() függvénnyel váltogatunk. Ezután a cudaMemcpyPeer() függvénnyel átmásoljuk a 2. videokártyán levő változó értékét a másik videokártyán levő változóba és ismét kiírjuk a két változó értékét.
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <stdio.h>
__global__ void ertek_beallitas(int *c, int szam)
{
c[0] = szam;
}
int main()
{
int *dev0_adat;
int *dev1_adat;
int akt_adat=0;
cudaSetDevice(0);
cudaMalloc((void**)&dev0_adat, 1 * sizeof(int));
cudaMemcpy(dev0_adat, &akt_adat, 1 * sizeof(int), cudaMemcpyHostToDevice);
cudaSetDevice(1);
cudaMalloc((void**)&dev1_adat, 1 * sizeof(int));
cudaMemcpy(dev0_adat, &akt_adat, 1 * sizeof(int), cudaMemcpyHostToDevice);
printf("Memoria lefoglalva...\n");
cudaSetDevice(0);
ertek_beallitas << <1, 1 >> > (dev0_adat,1);
cudaDeviceSynchronize();
cudaSetDevice(1);
ertek_beallitas << <1, 1 >> > (dev1_adat, 9);
cudaDeviceSynchronize();
printf("Kernelek ok...\n");
cudaSetDevice(0);
cudaMemcpy(&akt_adat, dev0_adat, 1 * sizeof(int), cudaMemcpyDeviceToHost);
printf("DEV0: %i\n", akt_adat);
cudaSetDevice(1);
cudaMemcpy(&akt_adat, dev1_adat, 1 * sizeof(int), cudaMemcpyDeviceToHost);
printf("DEV1: %i\n", akt_adat);
cudaMemcpyPeer(dev0_adat,0, dev1_adat,1,1*sizeof(int));
printf("Csere ok...\n");
cudaSetDevice(0);
cudaMemcpy(&akt_adat, dev0_adat, 1 * sizeof(int), cudaMemcpyDeviceToHost);
printf("DEV0: %i\n",akt_adat);
cudaSetDevice(1);
cudaMemcpy(&akt_adat, dev1_adat, 1 * sizeof(int), cudaMemcpyDeviceToHost);
printf("DEV1: %i\n", akt_adat);
cudaFree(dev0_adat);
cudaFree(dev1_adat);
return 0;
}
A program kimenete:
Memoria lefoglalva...
Kernelek ok...
DEV0: 1
DEV1: 9
Csere ok...
DEV0: 9
DEV1: 9
5.1.2 CUDA programok teljesítménymérése
CUDA programjaink viselkedéséről a legtöbb esetben rengeteget megtudhatunk a Windows Feladatkezelőjének ’Teljesítmény’ füle alatt található grafikonok, mérések tanulmányozásával is.
Amennyiben saját magunk is szeretnénk egyszerű méréseket végezni, ajánlható a GetTickCount, GetTickCount64, ill. a QueryPerformanceCounter Windows API függvények használata. Ezek bemutatásától eltekintük.
GPU monitorozása a Feladatkezelővel
A CUDA beépített mérőeszközökkel rendelkezik, melyek akkor igazán hasznosak, ha kódszinten vezérelt, egyedi teljesítménymérést szeretnénk végezni. A cudaEventCreate() és cudaEventRecord() függvénypáros segítségével pontosan meghatározható egy kernel futási ideje. Csak így tudjuk teljes bizonyossággal mérni a kerenelek futási idejét, mivel egyéb módszerek esetében elvileg az operációs rendszer által végzett egyéb műveletek is beleszámolódhatnak az időmérésbe. (A gyakorlatban nem eszik ennyire forrón a kását.)
Két időtartam különbségét ezredmásodpercben a cudaEventElapsedTime() függvénnyel tudjuk kiszámoltatni.
Az így létrehozott cudaEvent objektumokat minden esetben a cudaEventDestroy() függvénnyel törölnünk is kell, ha már nincsen rájuk szükség!
Egy példa a használatukra:
cudaEvent_t kezdes, befejezes;
cudaEventCreate(&kezdes);
cudaEventCreate(&befejezes);
cudaEventRecord(kezdes);
//******muveletek**********
cudaEventRecord(befejezes);
cudaEventSynchronize(befejezes);
float idotartam = 0; //***ez lesz a meres erteke***
cudaEventElapsedTime(&idotartam, kezdes, befejezes);
cudaEventDestroy(&kezdes);
cudaEventDestroy(&befejezes);
5.1.3 Futtatási profilok elemzése
Lehetőség van a CUDA programok futását automatikus teljesítménymérésnek is alávetni és ehhez még külön programoznunk sem kell. Az NVIDIA ehhez az Nsight Systems és az Nsight Compute eszközöket biztosítja. Az ide tartozó Visual Profiler segédprogramot (NVVP.EXE) az alábbi helyen találjuk:
C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA\v10.0\libnvvp
Az első indítás során meg kell adnunk az ún. Workspace elérési útját. Ide mentődnek el az ún. elemzési profilok.
A File – New Session menüpontot kiválasztva a ’’Create New Session” párbeszédablak jelenik meg, ahol a ’File’ mellett meg kell adnunk a lefordított CUDA programunk elérési útját.
A Next, majd a Finish gombra kattintva elindul alkalmazásunk, amit az egyébként megszokott módon használhatunk. Elegendő lehet pár másodpercig futtatni az alkalmazást, az alkalmazás bezárása után megjelennek a mérési adatok. Csak arra figyeljünk, hogy CUDA programrész biztosan végrehajtódjon.
Az alábbi példában egy 2 videokártyát párhuzamos renderelésre használó program elemzési adatai láthatóak.
Megjegyzendő, hogy az NVIDIA FrameView néven egy önálló monitorozó segédprogramot is kibocsátott, melyet szintén teljesítménymérésre használhatunk.
Érdemes megnézni az NVIDIA által összeállított „kihasználtsági” kalkulátort, mely statikusan képes kiszámítani és megjeleníteni az adott architektúrára levetítve a GPU erőforrásainak kihasználtságát. Ez egy interaktív Excel táblázat, mely az alábbi oldalról tölthető le:
https://docs.nvidia.com/cuda/cuda-occupancy-calculator/index.html
6 Függelék
6.1 CUDA alapú renderlő algoritmusok
A nagyteljesítményű adatfeldolgozást praktikus lehet valamilyen grafikus megjelenítéssel demonstrálni. Erre áll itt egy példa. A Pegazus renderlő motorom ugyanis CUDA kompatibilis is. A Pegazus egyebek mellett Direct2D-t is képes használni a renderelt képtartalmak megjelenítéséhez, de a képkockák teljes mértékben a CUDA eszközökön vannak előállítva, képponttömbökben, kissé szakmaibb nyelven megfogalmazva: frame bufferekben. A renderelés mindössze tömbelemek állítgatásából áll.
A témát itt csak érintjük, de akit rélszleteiben is érdekel, annak ajánlhatóak az alábbi leírások is:
• 3D programozás
• Szupergyors rajzoló algoritmusok
• Direct2D programozás dióhéjban.
A Pegazus az alakzatok leíró információit tömbökben tárolja. A tömbökből egy kernel szedegeti ki az alakzatok csúcspontjainak koordinátáit, ez a kernel lesz párhuzamosan végrehajtva. Egy elemi alakzatot egy CUDA mag egy nekifutásra rajzol ki, tehát ez a lépés már nincsen tovább bontva, mivel az ugrásszerűen megnövelné a kernelindítások számát, ami viszont már komoly teljesítménycsökkenést vonna maga után.
Ezeket az algoritmusokat találjuk meg itt, sőt egy egyszerű teljesítménymérő alkalmazás forráskódját is.
6.1.1 Pufferelés megvalósítása
A legegyszerűbb esetben két képpuffert alkalmazunk: egyet a CUDA eszközön állítunk elő, a renderelés során. A másodikat a CPU kezeli és csak formális: ide másoljuk a CUDA eszközön renderelt képpuffert, a tényleges megjelenítéshez. A host kód szemszögéből nézve mindig csak egy képponttömböt kell kirajzolni.
A képpuffer törlését CUDA-val az alábbi függvény végzi:
void CUDA_cleanup_main_buffer(void)
{
cudaMemset(dev_kepadat, 255, KEPERNYO_HEIGHT * KEPERNYO_WIDTH * sizeof(unsigned int));
}
Az alábbi kód rajzolja ki a képpuffert a képernyőre:
void swap_main_buffer(void)
{
hbitmapScreen = CreateBitmap(SCREEN_WIDTH, SCREEN_HEIGHT, 1, 8 * 4, (void*)memkeptarolo);
SelectObject(hdcMemDC, hbitmapScreen);
BitBlt(hdcWindow, 0, 0, SCREEN_WIDTH, SCREEN_HEIGHT, hdcMemDC, 0, 0, SRCCOPY);
DeleteObject(hbitmapScreen);
}
6.1.2 Elemi alakzatok kirajzolása CUDA kernelekkel
Ponttömb kirajzolását az alábbi kernellel végezhetjük el. A vonal- és háromszögrajzoláshoz használt CUDA_DrawLine és CUDA_FillTriangle függvények kódjai megtalálhatóak a leírás végi példaalkalmazás teljes forráskódjában.
__device__ void CUDA_ SetPixel_main_buffer(int x1, int y1, int szin, unsigned int *puffer)
{
puffer[(y1 * KEPERNYO_WIDTH) + x1] = szin;
}
A példaprogramunkban a D2D_rajzolas() függvény végzi el egy képkocka előállítását. A render_objects() függvény a belépési pontja a CUDA alapú renderelésnek.
void D2D_rajzolas(ID2D1HwndRenderTarget* pRT)
{
CUDA_cleanup_main_buffer();
render_objects<<<blokkok,szalak >>>(dev_kepadat,dev_zbuffer);
cudaDeviceSynchronize();
cudaMemcpy(kepadat, dev_kepadat, KEPERNYO_WIDTH * KEPERNYO_HEIGHT * sizeof(unsigned int), cudaMemcpyDeviceToHost);
swap_main_buffer();
}
Alább található egy komplex példaprogram (CUDA_09.CU) teljes forráskódja, ami ezeket az algoritmusokat használja.
Néhány kiegészítés és magyarázat a program működéséhez.
A render_objects függvény minimális módosításával befolyásolhatjuk, hogy pontokat, vonalakat, vagy háromszögeket rajzoljon a program. Alább bemutatjuk mindhárom változatot, vastagon kiemelve a különbségeket tartalmazó sorokat.
Pontrajzolás
__global__ void render_objects(int maxitemcount, float *arrayX, float *arrayY, unsigned int *colorpuffer, unsigned int *puffer)
{
int i, px, py, tesztcolor;
int index = (blockIdx.x * blockDim.x) + (threadIdx.x * 1);
int stride = blockDim.x * gridDim.x;
for (i = index; i < maxitemcount; i += stride)
{
CUDA_SetPixel(arrayX[i], arrayY[i], colorpuffer[i], puffer);
}
}
Vonalrajzolás
__global__ void render_objects(int maxitemcount, float *arrayX, float *arrayY, unsigned int *colorpuffer, unsigned int *puffer)
{
int i, px, py, tesztcolor;
int index = (blockIdx.x * blockDim.x) + (threadIdx.x * 2);
int stride = blockDim.x * gridDim.x;
for (i = index; i < maxitemcount - 1; i += stride)
{
CUDA_DrawLine(arrayX[i], arrayY[i], arrayX[i + 1], arrayY[i + 1],
colorpuffer[i], puffer);
}
}
Háromszögrajzolás
__global__ void render_objects(int maxitemcount, float *arrayX, float *arrayY, unsigned int *colorpuffer, unsigned int *puffer)
{
int i, px, py, tesztcolor;
int index = (blockIdx.x * blockDim.x) + (threadIdx.x * 3);
int stride = blockDim.x * gridDim.x;
for (i = index; i < maxitemcount - 2; i += stride)
{
CUDA_FillTriangle(arrayX[i], arrayY[i], arrayX[i + 1], arrayY[i + 1],
arrayX[i + 2], arrayY[i + 2], colorpuffer[i], puffer);
}
}
A program statisztikát is tárol a STATISZTIKA.TXT fájlban, minden egyes képkocka kirajzolásához.
Példa egy ilyen mérési eredményre:
Képpuffer törlése: : 0
Rendereléshez szükséges idő: : 906
Képkocka másolása és megjelenítése: : 16
Képpuffer törlése: : 0
Rendereléshez szükséges idő: : 31
Képkocka másolása és megjelenítése: : 15
Amennyiben CUDA hiba keletkezne, az kiíródik a programablak bal felső sarkába.
A program tetszés szerint továbbfejleszthető általános grafikus megjelenítést használó alkalmazássá, de akár teljesítménymérési célokra is hasznosítani lehet.
A teljes forráskód:
#include "cuda_runtime.h"
#include "device_launch_parameters.h"
#include <math.h>
#include <stdio.h>
#include <windows.h>
#include <time.h>
#include <d2d1.h>
#include <d2d1helper.h>
#pragma comment(lib, "d2d1")
//*****double buffering*****
#define KEPERNYO_WIDTH 600
#define KEPERNYO_HEIGHT 400
D2D1_RECT_U display_area;
ID2D1Bitmap *memkeptarolo = NULL;
unsigned int kepadat[KEPERNYO_WIDTH * KEPERNYO_HEIGHT];
//**************************************
//**********STATISZTIKA*******************
int kezdet, vege;
FILE *statfajl;
void writestat(char *szoveg, int ertek);
void meres_start(void);
void meres_end(void);
int getrandom(int maxnum);
//***************************************
//**************PEGAZUS************
#define MAX_OBJ_NUM 1000000
float raw_verticesX[MAX_OBJ_NUM], raw_verticesY[MAX_OBJ_NUM];
int raw_colors[MAX_OBJ_NUM];
int raw_vertices_length;
//*******CUDA*************
float *dev_raw_verticesX, *dev_raw_verticesY;
unsigned int *dev_raw_colors;
unsigned int *dev_kepadat;
//************************
void data_transfer_to_GPU(void);
void D2D_rajzolas(ID2D1HwndRenderTarget* pRT);
__global__ void render_objects(int maxitemcount, float *arrayX, float *arrayY, unsigned int *colorpuffer, unsigned int *puffer);
//************************************
//***********STANDARD WIN32API WINDOWING************
ID2D1Factory* pD2DFactory = NULL;
ID2D1HwndRenderTarget* pRT = NULL;
#define HIBA_00 TEXT("Error:Program initialisation process.")
HINSTANCE hInstGlob;
int SajatiCmdShow;
char szClassName[] = "WindowsApp";
HWND Form1; //Ablak kezeloje
LRESULT CALLBACK WndProc0(HWND, UINT, WPARAM, LPARAM);
//******************************************************
//*****double buffering*****
void create_main_buffer(void);
void CUDA_cleanup_main_buffer(void);
void swap_main_buffer(void);
//**************************************
//*****drawig algorithms*****
__device__ void CUDA_SetPixel(int x1, int y1, int color, unsigned int *puffer);
__device__ void CUDA_DrawLine(int x1, int y1, int x2, int y2, int color, unsigned int *puffer);
__device__ void CUDA_FillTriangle(int x1, int y1, int x2, int y2, int x3, int y3, int color, unsigned int *puffer);
//**************************************
//*********************************
//The main entry point of our program
//*********************************
int WINAPI WinMain(HINSTANCE hInstance, HINSTANCE hPrevInstance, PSTR szCmdLine, int iCmdShow)
{
static TCHAR szAppName[] = TEXT("StdWinClassName");
HWND hwnd;
MSG msg;
WNDCLASS wndclass0;
SajatiCmdShow = iCmdShow;
hInstGlob = hInstance;
//*********************************
//Preparing Windows class
//*********************************
wndclass0.style = CS_HREDRAW | CS_VREDRAW;
wndclass0.lpfnWndProc = WndProc0;
wndclass0.cbClsExtra = 0;
wndclass0.cbWndExtra = 0;
wndclass0.hInstance = hInstance;
wndclass0.hIcon = LoadIcon(NULL, IDI_APPLICATION);
wndclass0.hCursor = LoadCursor(NULL, IDC_ARROW);
wndclass0.hbrBackground = (HBRUSH)GetStockObject(LTGRAY_BRUSH);
wndclass0.lpszMenuName = NULL;
wndclass0.lpszClassName = TEXT("WIN0");
//*********************************
//Registering our windows class
//*********************************
if (!RegisterClass(&wndclass0))
{
MessageBox(NULL, HIBA_00, TEXT("Program Start"), MB_ICONERROR);
return 0;
}
//*********************************
//Creating the window
//*********************************
Form1 = CreateWindow(TEXT("WIN0"),
TEXT("CUDA - DIRECT2D"),
(WS_OVERLAPPED | WS_SYSMENU | WS_THICKFRAME | WS_MAXIMIZEBOX | WS_MINIMIZEBOX),
50,
50,
KEPERNYO_WIDTH,
KEPERNYO_HEIGHT,
NULL,
NULL,
hInstance,
NULL);
//*********************************
//Displaying the window
//*********************************
ShowWindow(Form1, SajatiCmdShow);
UpdateWindow(Form1);
//*********************************
//Activating the message processing for our window
//*********************************
while (GetMessage(&msg, NULL, 0, 0))
{
TranslateMessage(&msg);
DispatchMessage(&msg);
}
return msg.wParam;
}
//*********************************
//The window's callback funtcion: handling events
//*********************************
LRESULT CALLBACK WndProc0(HWND hwnd, UINT message, WPARAM wParam, LPARAM lParam)
{
HDC hdc;
PAINTSTRUCT ps;
switch (message)
{
//*********************************
//When creating the window
//*********************************
case WM_CREATE:
srand((unsigned)time(NULL));
D2D1CreateFactory(D2D1_FACTORY_TYPE_SINGLE_THREADED, &pD2DFactory);
pD2DFactory->CreateHwndRenderTarget(
D2D1::RenderTargetProperties(),
D2D1::HwndRenderTargetProperties(
hwnd, D2D1::SizeU(KEPERNYO_WIDTH, KEPERNYO_HEIGHT)),
&pRT);
cudaDeviceReset();
create_main_buffer();
cudaMalloc((void**)&dev_raw_verticesX, MAX_OBJ_NUM * sizeof(float));
cudaMalloc((void**)&dev_raw_verticesY, MAX_OBJ_NUM * sizeof(float));
cudaMalloc((void**)&dev_raw_colors, MAX_OBJ_NUM * sizeof(unsigned int));
cudaMalloc((void**)&dev_kepadat, KEPERNYO_WIDTH * KEPERNYO_HEIGHT * sizeof(unsigned int));
int i;
for (i = raw_vertices_length = 0; i < MAX_OBJ_NUM; ++i)
{
raw_verticesX[i] = getrandom(KEPERNYO_WIDTH);
raw_verticesY[i] = getrandom(KEPERNYO_HEIGHT);
raw_colors[i] = RGB(getrandom(255), getrandom(255), getrandom(255));
++raw_vertices_length;
}
data_transfer_to_GPU();
return 0;
//*********************************
//to eliminate color flickering
//*********************************
case WM_ERASEBKGND:
return (LRESULT)1;
//*********************************
//Repainting the client area of the window
//*********************************
case WM_PAINT:
hdc = BeginPaint(hwnd, &ps);
EndPaint(hwnd, &ps);
D2D_rajzolas(pRT);
return 0;
//*********************************
//Closing the window, freeing resources
//*********************************
case WM_CLOSE:
pRT->Release();
pD2DFactory->Release();
cudaFree(dev_raw_verticesX);
cudaFree(dev_raw_verticesY);
cudaFree(dev_raw_colors);
cudaFree(dev_kepadat);
DestroyWindow(hwnd);
return 0;
//*********************************
//Destroying the window
//*********************************
case WM_DESTROY:
PostQuitMessage(0);
return 0;
}
return DefWindowProc(hwnd, message, wParam, lParam);
}
//********************************
//PEGAZUS
//********************************
void create_main_buffer(void)
{
pRT->CreateBitmap(D2D1::SizeU(KEPERNYO_WIDTH, KEPERNYO_HEIGHT),
D2D1::BitmapProperties(D2D1::PixelFormat(DXGI_FORMAT_B8G8R8A8_UNORM,
D2D1_ALPHA_MODE_IGNORE)), &memkeptarolo);
}
void CUDA_cleanup_main_buffer(void)
{
cudaMemset(dev_kepadat, 200, KEPERNYO_HEIGHT*KEPERNYO_WIDTH * sizeof(unsigned int));
}
void swap_main_buffer(void)
{
display_area.left = 0;
display_area.top = 0;
display_area.right = KEPERNYO_WIDTH;
display_area.bottom = KEPERNYO_HEIGHT;
memkeptarolo->CopyFromMemory(&display_area, kepadat, KEPERNYO_WIDTH * sizeof(unsigned int));
pRT->BeginDraw();
pRT->DrawBitmap(memkeptarolo, D2D1::RectF(0.0f, 0.0f, KEPERNYO_WIDTH, KEPERNYO_HEIGHT), 1.0f, D2D1_BITMAP_INTERPOLATION_MODE_NEAREST_NEIGHBOR, NULL);
pRT->EndDraw();
}
__device__ void CUDA_SetPixel(int x1, int y1, int color, unsigned int *puffer)
{
puffer[(y1 * KEPERNYO_WIDTH) + x1] = color;
}
__device__ void CUDA_DrawLine(int x1, int y1, int x2, int y2, int color, unsigned int *puffer)
{
bool flip = false;
int swap, offset;
if (abs(x2 - x1) < 2 && abs(y2 - y1) < 2)
{
puffer[(y2*KEPERNYO_WIDTH) + x2] = color; return;
}
if (abs(x1 - x2) < abs(y1 - y2))
{
swap = x1;
x1 = y1;
y1 = swap;
swap = x2;
x2 = y2;
y2 = swap;
flip = true;
}
if (x1 > x2)
{
swap = x1;
x1 = x2;
x2 = swap;
swap = y1;
y1 = y2;
y2 = swap;
}
int dx = x2 - x1;
int dy = y2 - y1;
int marker1 = abs(dy) * 2;
int marker2 = 0;
int y = y1, x;
if (flip)
{
for (x = x1; x <= x2; ++x)
{
offset = (x * KEPERNYO_WIDTH);
puffer[offset + y] = color;
marker2 += marker1;
if (marker2 > dx)
{
y += (y2 > y1 ? 1 : -1);
marker2 -= dx * 2;
}
}
}
else
{
for (x = x1; x <= x2; ++x)
{
offset = (y * KEPERNYO_WIDTH);
puffer[offset + x] = color;
marker2 += marker1;
if (marker2 > dx)
{
y += (y2 > y1 ? 1 : -1);
marker2 -= dx * 2;
}
}
}
}
__device__ void CUDA_FillTriangle(int x1, int y1, int x2, int y2, int x3, int y3, int color, unsigned int *puffer)
{
int Ax, Ay, Bx, By, i, j;
int swapx, swapy, offset, maxoffset = KEPERNYO_HEIGHT * KEPERNYO_WIDTH;
if (y1 == y2 && y1 == y3) return;
if (y1 > y2)
{
swapx = x1;
swapy = y1;
x1 = x2;
y1 = y2;
x2 = swapx;
y2 = swapy;
}
if (y1 > y3)
{
swapx = x1;
swapy = y1;
x1 = x3;
y1 = y3;
x3 = swapx;
y3 = swapy;
}
if (y2 > y3)
{
swapx = x3;
swapy = y3;
x3 = x2;
y3 = y2;
x2 = swapx;
y2 = swapy;
}
int t_height = y3 - y1;
for (i = 0; i < t_height; ++i)
{
bool lower_part = i > y2 - y1 || y2 == y1;
int part_height = lower_part ? y3 - y2 : y2 - y1;
float alpha = (float)i / t_height;
float beta = (float)(i - (lower_part ? y2 - y1 : 0)) / part_height;
Ax = x1 + (x3 - x1)*alpha;
Ay = y1 + (y3 - y1)*alpha;
Bx = lower_part ? x2 + (x3 - x2)*beta : x1 + (x2 - x1)*beta;
By = lower_part ? y2 + (y3 - y2)*beta : y1 + (y2 - y1)*beta;
if (Ax > Bx)
{
swapx = Ax;
swapy = Ay;
Ax = Bx;
Ay = By;
Bx = swapx;
By = swapy;
}
offset = (y1 + i)*KEPERNYO_WIDTH;
for (j = Ax; j < Bx; ++j)
{
if (offset + j > maxoffset) continue;
puffer[offset + j] = color;
}
}
}
void data_transfer_to_GPU(void)
{
cudaMemcpy(dev_raw_verticesX, raw_verticesX, raw_vertices_length * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dev_raw_verticesY, raw_verticesY, raw_vertices_length * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(dev_raw_colors, raw_colors, raw_vertices_length * sizeof(unsigned int), cudaMemcpyHostToDevice);
}
void D2D_rajzolas(ID2D1HwndRenderTarget* pRT)
{
char hibauzenet[256];
meres_start();
CUDA_cleanup_main_buffer();
cudaDeviceSynchronize();//opcionálisan elhagyható
meres_end();
writestat("Képpuffer törlése: ", vege);
strcpy_s(hibauzenet, cudaGetErrorString(cudaGetLastError()));
SetWindowTextA(Form1, hibauzenet);
meres_start();
int szalak_szama = 128;
int blokkok_szama = (100000 + szalak_szama - 1) / szalak_szama;
render_objects << <80,8 >> > (raw_vertices_length, dev_raw_verticesX, dev_raw_verticesY, dev_raw_colors, dev_kepadat);
cudaDeviceSynchronize();
meres_end();
writestat("Rendereléshez szükséges idő: ", vege);
strcpy_s(hibauzenet, cudaGetErrorString(cudaGetLastError()));
SetWindowTextA(Form1, hibauzenet);
meres_start();
cudaMemcpy(kepadat, dev_kepadat, KEPERNYO_WIDTH * KEPERNYO_HEIGHT * sizeof(unsigned int), cudaMemcpyDeviceToHost);
strcpy_s(hibauzenet, cudaGetErrorString(cudaGetLastError()));
SetWindowTextA(Form1, hibauzenet);
swap_main_buffer();
meres_end();
writestat("Képkocka másolása és megjelenítése: ", vege);
}
__global__ void render_objects(int maxitemcount, float *arrayX, float *arrayY, unsigned int *colorpuffer, unsigned int *puffer)
{
int i, px, py, tesztcolor;
int index = (blockIdx.x * blockDim.x) + (threadIdx.x * 2);
int stride = blockDim.x * gridDim.x;
for (i = index; i < maxitemcount - 1; i += stride)
{
//CUDA_SetPixel(arrayX[i], arrayY[i], colorpuffer[i], puffer);
CUDA_DrawLine(arrayX[i], arrayY[i], arrayX[i + 1], arrayY[i + 1], colorpuffer[i], puffer);
//CUDA_FillTriangle(arrayX[i], arrayY[i], arrayX[i + 1], arrayY[i + 1], arrayX[i + 2], arrayY[i + 2], colorpuffer[i], puffer);
}
}
void meres_start(void)
{
kezdet = GetTickCount();
}
void meres_end(void)
{
vege = GetTickCount() - kezdet;
}
void writestat(char *szoveg, int ertek)
{
statfajl = fopen("statisztika.txt", "at");
if (statfajl == NULL) return;
fprintf(statfajl, "%s: ", szoveg);
fprintf(statfajl, "%i\n", ertek);
fclose(statfajl);
}
int getrandom(int maxnum)
{
return (double)rand() / (RAND_MAX + 1) * maxnum;
}
6.2 A leírásban használt CUDA függvények
Általános eszközkezelés
__host__ cudaError_t cudaDeviceReset ( void )
__host__ __device__ cudaError_t cudaDeviceSynchronize ( void )
__host__ __device__ cudaError_t cudaGetDeviceCount ( int* count )
__host__ cudaError_t cudaGetDeviceProperties ( cudaDeviceProp* prop, int device )
__host__ cudaError_t cudaSetDevice ( int device )
Memóriakezelés
__host__ cudaError_t cudaMallocManaged ( void** devPtr, size_t size, unsigned int flags = cudaMemAttachGlobal )
__host__ __device__ cudaError_t cudaFree ( void* devPtr )
__host__ __device__ cudaError_t cudaMalloc ( void** devPtr, size_t size )
__host__ cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count, cudaMemcpyKind kind )
__host__ cudaError_t cudaMemset ( void* devPtr, int value, size_t count )
__host__ cudaError_t cudaMemcpyPeer ( void* dst, int dstDevice, const void* src, int srcDevice, size_t count )
Teljesítménymérés
__host__ cudaError_t cudaEventCreate ( cudaEvent_t* event )
__host__ __device__ cudaError_t cudaEventRecord ( cudaEvent_t event, cudaStream_t stream = 0 )
__host__ cudaError_t cudaEventSynchronize ( cudaEvent_t event )
__host__ __device__ cudaError_t cudaEventDestroy ( cudaEvent_t event )
__host__ cudaError_t cudaEventElapsedTime ( float* ms, cudaEvent_t start, cudaEvent_t end )
Hibakezelés
__host__ __device__ const char* cudaGetErrorString ( cudaError_t error )
__host__ __device__ cudaError_t cudaGetLastError ( void )