Základy zpracování signálů na GPU

Rok / Year: 2010

Svazek / Volume: 12

Číslo / Number: 2

Základy zpracování signálů na GPU GPU signal processing basics Ondřej Morský [email protected] Fakulta elektrotechniky a komunikačních technologií VUT v Brně

Abstrakt: Článek seznamuje čtenáře s možnostmi moderních grafických karet výrobce nVidia pro použití při náročných matematických výpočtech. Dále je zde popsán jednouchý způsob zpracování signálu (korelace) pomocí dnes nejpoužívanějších programovacích jazyků C++, C# a SSE a C for CUDA.

Abstract: This paper describes possibilities of modern graphic cards manufactured by nVidia for using in difficult mathematical operations. In next part, there is description of simple signal processing (correlation) using today’s most popular programming languages and technologies C++, C# and SSE and C for CUDA.

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

ZÁKLADY ZPRACOVÁNÍ SIGNÁLŮ NA GPU Ing. Ondřej Morský Email: [email protected] Článek seznamuje čtenáře s možnostmi moderních grafických karet výrobce nVidia pro použití při náročných matematických výpočtech. Dále je zde popsán jednouchý způsob zpracování signálu (korelace) pomocí dnes nejpoužívanějších programovacích jazyků – C++, C# a SSE a C for CUDA.

grafických karet výrazně rychlejší s širší sběrnicí než běžná RAM.

1. ÚVOD V dnešní době výpočetní výkon hardwaru neustále roste. Platí to nejen pro procesory, ale i pro grafické karty. Ty byly donedávna vývojáři softwaru přehlíženy. A to často pouze proto, že jde o zařízení určené pro hry a seriózní programátoři se jím tedy nezabývají. I přesto se podařilo firmám zabývajícím se grafikou a hrami prosadit tento hardware jako alternativu při náročných výpočtech k moderním CPU1.

2.2.1. T RANSFORM AND LIGHTNING V roce 1999 se objevil první 3D grafický akcelerátor – nVidia GeForce 256. [2] Z dnešního pohledu prakticky nic neuměl. Jednalo se o procesor, který byl určen pro násobení čtyř-rozměných vektorů a matic reálných čísel. Prováděl geometrické transformace zadaných bodů pomocí nastavených matic a následně rasterizaci a texturování. Zpracování dat bylo možné ovlivnit pouze tím, že se jednotlivé bloky (např. výpočet osvětlení) vypnuly, nebo zapnuly, případně se jim nastavily parametry (směr světla, …). Jednotka vykonávající tyto operace byla označována jako T&L – Transform (clipping) and lightning. GeForce 256 měla čtyři tyto jednotky zapojené paralelně.

Tento článek je zaměřen na běžné domácí počítače platformy Windows x86/x64 a grafické karty nVidia.

2. HISTORIE V této kapitole bude porovnán vývoj CPU a GPU 2a zároveň vypsány nejdůležitější body ve vývoji grafických procesorů, které směřovaly k použití GPU pro všeobecné výpočty.

2.2.2. S HADERY V roce 2001 se na trhu objevil grafický akcelerátor označený jako GeForce 3, který umožnil jednotku T&L vypnout a její funkci nahradit programovatelnou jednotkou označovanou jako shader. Přestože programy pro tyto jednotky byly velmi omezené – především počtem instrukcí, způsobem práce s pamětí apod., udala tato technologie směr dnešního vývoje.

2.1. CPU Historie CPU je všeobecné známá a není nutné ji zde podrobně popisovat. Důležité je, že CPU byl od svého vzniku určen k všeobecnému použití – je možné ho naprogramovat. Z toho vyplývá jeho univerzálnost. Je možné ho tedy velmi jednoduše použít pro jakýkoliv typ výpočtu. Při jeho vývoji byl kladen důraz na výkon – především frekvenci a případně rozšiřující instrukce. K výrazné změně došlo až v dubnu roku 2005, kdy Intel začal prodávat svůj první více-jádrový procesor. [1] Od té doby je vidět změna ve vývoji CPU směrem k většímu paralelismu při zachování stejných frekvencí.

GPU obsahovaly dva shadery, které byly realizovány různými částmi čipu. První typ označovaný3 jako Vertex Shader byl určen pro zpracování geometrický dat – tj. pro násobení matic a vektorů. Druhý typ – Pixel Shader byl určen pro výpočty jednotlivých pixelů výsledného obrazu – tj. práce s texturami a barvami. GeForce 3 měla stále pouze 4 jednotky.

2.2. GPU

2.2.3. U NIFIED S HADER MODEL

Grafické procesory byly na rozdíl od CPU vždy určeny pouze pro výpočty vhodné pro vytváření a zpracování obrazových dat a to především pro operace s čísly s plovoucí desetinnoou čárkou. Důležité je, že již od začátku jsou navrhovány pro paralelní výpočty. Často používanou funkcí je také hardwarová podpora některých filtračních a komprimačních algoritmů. Další podstatný rozdíl oproti CPU spočívá v tom, že jsou určeny pro práci s objemnými obrazovými daty, proto je i paměť

Společně s DirectX 104 přišla další podstatná změna. Přibyl nový typ – Geometry Shader, který umožňuje generovat nové 3D objekty přímo na GPU, Shadery dospěly do čtvrté verze označované jako Shader Model 4.0, nebo také jako Unified Shader Model. Jak již z názvu vyplývá, jedná se o sloučení všech tří typů do jednoho. Znamená to, že vše je realizováno jednou částí

1

CPU – Central processing unit

3

Názvosloví Direct3D

2

GPU – Graphics processing unit

4

rok 2006

13-1

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

není možné, aby vlastnosti bodu závisely na jiných bodech, nebo na jiných výsledcích → nejsou potřeba žádné synchronizační funkce. Je proto možné plně využít paralelismus grafické karty – až desítky4 bodů mohou být zpracovány současně.

integrovaného obvodu a především, že ve všech typech je možné použít stejné instrukce. Nové univerzální jednotky jsou nazývány jako Stream procesory. Jejich počet na GeForce 8 se pohybuje, podle verze, v rozmezí 8 až 128.

2.3. GPGPU1 General-purpose computing on graphics processing unit je technika používání GPU pro výpočty, které obvykle vykonává CPU. Jedná se o několik úprav GPU pro zvýšení přesnosti výpočtu a zjednodušení programování. [3]

3.

GPU sestaví z výsledných 2D bodů trojúhelníky a provede rasterizaci – označí pixely obrazovky patřící trojúhelníkům.

4. Pro každý takovýto pixel je spuštěn Pixel shader. Ten obvykle přečte barvu z textury uložené v paměti grafické karty a označí touto barvou pixel. Zde je možné opět využít paralelismus a spočítat několik desítek pixelů současně. Opět je nutné, aby počet vstupních pixelů odpovídal počtu výstupních.

2.3.1. B ROOK GPU BrookGPU je první programové prostředí, které umožnilo snadno využít grafické procesory k výpočtům. Využívalo k tomu Direct3D, OpenGL, nebo Close To Metal.

2.3.2. CUDA

Hlavní nevýhody shaderů oproti CPU

CUDA2 je rozšíření běžných stream procesorů o funkce, které jsou nutné pro matematické výpočty. Jde o produkt firmy nVidia fungující pod Win32/64, Linux 32/64 a Mac OS. CUDA je založeno na BrookGPU.

2.3.3. D IRECT X 11 C OMPUTE Hlavní nevýhoda CUDA spočívá v tom, že je určená pouze pro produkty firmy nVidia a to řady GeForce 8 a vyšší. Součástí příští verze DirectX (Microsoft) bude rozšíření označené Compute, které by mělo sjednotit používání grafických karet pro všeobecné výpočty od všech výrobců do jednotného programátorského přístupu. Nevýhodou bude, že je určen pouze pro OS Windows a je nutný hardware podporující DirectX 11.

3. CUDA

•

Je možné pracovat pouze s datovými typy float. GPU neumí celočíselné typy, a tedy ani bitové operace.

•

Počet vstupů je vždy stejný jako počet výstupů. Výsledek 10. vstupního bodu bude vždy na 10. výstupním místě.

•

Vždy je nutné spustit minimálně dva programy – Vertex a Pixel. Geometry shader je volitelný.

•

Data v průběhu zpracovávání je možné číst pouze z textury – pomocí normalizovaných souřadnic.

•

Není možná synchronizace.

Hlavní výhody shaderů oproti CPU

Pro pochopení, k čemu vlastně CUDA slouží, je nutné znát alespoň základy3 vytváření 3D obrazu na PC:

•

Když běží GPU, je CPU nezatížen, a může tedy vykonávat jinou činnost

1. CPU vytvoří pole 3D bodů, které vždy po třech představují jeden trojúhelník v prostoru, a uloží ho do paměti grafické karty

•

Extrémní rychlost matematických výpočtů díky paralelním jednotkám pracujících na vysokých frekvencích.

2. Body jsou zpracovány pomocí geometrických transformací ze 3D souřadnic v prostoru do 2D souřadnic na monitoru. GPU začne číst paměť tak, že přečte JEDEN bod a předá ho jednomu stream procesoru (Vertex shaderu), a očekává vždy JEDEN výsledek. Není možné, aby stream procesor bod přidal, nebo odstranil. Pořadí zpracování bodů může být náhodné (Pořadí vstupních a výsledných bodů je však vždy stejné). Počet bodů na vstupu je vždy stejný jako počet bodů na výstupu. Zde je důležité upozornit, že

1

Někdy GPGP, nebo pouze GP

2

Compute unified device architecture

Zjednodušeně lze říci, že na GPU lze naprogramovat pouze něco, ale když to jde, je to i bez jakýchkoliv optimalizací rychlejší než na CPU. CUDA je, jak již bylo řečeno, rozšíření stream procesorů o funkce, které umožní širší využití grafických karet. Na internetu koluje spousta informací o tom, že CUDA je jazyk. Toto tvrzení je chybné. Celý název je totiž CUDA Architecture, což je společně s hardwarovými úpravami grafického procesoru také několik softwarových knihoven a rozšíření do nejčastějších programovacích jazyků – nejznámější C for CUDA.

3

4

Velmi zjednodušený příklad pro použití pouze Vertex a Pixel shaderů

V závislosti na GPU se může jednat až o stovky, v případě použití více GPU/grafických karet až tisíce.

13-2

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

3.2.3. K ERNEL Nejdůležitější změny se týkají právě hardwaru. •

Kernel je funkce běžící v každém vláknu na GPU. Jde vlastně o program, který GPU vykoná.

Je možné pracovat s celočíselnými datovými typy a používat bitové operace

•

Pole pro výstup musí být připraveno předem, ale výsledek nemusí být uložen

•

Výsledek může přepsat vstupní data.

•

Spouští se pouze jeden program.

•

Data je možné číst z několika typů pamětí, přičemž textura s normalizovanými souřadnicemi je pouze jedna z možností.

__global__ void funkce(float *A,float *B) { … }

3.2.4. T HREAD

•

Je možné synchronizovat skupiny threadů – tzv. bloky. (viz. Názvosloví)

•

Několik atomických1 funkcí

•

Pro programátora není nutné znát grafické API, ale přesto je možné kombinovat CUDA a Direct3D, nebo OpenGL

Thread je běžící kernel – podobně jako na běžném CPU.

3.2.5. B LOCK Block je skupina threadů. Ty jsou do bloku vkládány jako do dvou až trojrozměrného pole. To umožňuje snazší zpracování tří-rozměrných dat – každý thread zpracuje prvek na své pozici. Velikost bloku je omezena3 celkovým počtem threadů. Vlákna v jednom bloku mají přístup ke společné paměti – Shared Memory a je možné je vzájemně synchronizovat.

běžících

3.1. K ČEMU VLASTNĚ?

3.2.6. G RID

GPGPU se používá nejčastěji k vědeckým účelům. na stránkách www.nvidia.com/cuda nebo www.amd.com/stream je několik příkladů, kde jsou tyto technologie využívány v praxi. Velmi časté je použití pro simulace počasí, nebo pro analýzu dat zachycených teleskopy. Poměrně známý program pro distribuované vědecké výpočty BOINC2 existuje také ve verzi využívající právě CUDA - www.gpugrid.net.

Grid je až dvourozměrná skupina bloků. Vlákna ani bloky nemohou sdílet data uvnitř gridu. Každý thread může zjistit, jak velký je grid i block a na jakých souřadnicích dané struktury běží.

Thread

3.2. NÁZVOSLOVÍ Pojmenování jednotlivých částí je velmi složité. Zde bude popsána pouze nutná část. Uváděné příklady jsou napsány v C for CUDA. Více podrobností a příkladů je možné nalézt zde [4].

Block

3.2.1. H OST Z pohledu programování grafické karty je CPU nadřízený procesor, a říká se mu proto host.

3.2.2. D EVICE GPU je zařízení, na kterém běží výpočty, proto se označuje jako device. Grid

1

Funkce, u kterých je zajištěno, že žádné jiné vlákno nepřitoupí ke stejné části paměti během provádění této funkce. Např: atomicAdd() přečte data z paměti, přičte číslo a uloží zpět

Obr. č. 1: Hierarchie vláken

2

Software pro distribuované výpočty. Vstupní data rozdělí po síti mezi několik PC a poté posbírá výsledky.

3

13-3

na GeForce 8600 je to 512 threadů

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

3.3. ROZDĚLENÍ PAMĚTI

4.1. IMPLEMENTACE

Podrobnější informace je opět možné nalézt zde [4].

Pro změření výkonu byly zvoleny dvě technologie vhodné pro tento typ výpočtu.

3.3.1. R EGISTERS

•

SSE - pomocí C++ a assembleru

Jde o nejrychlejší paměť. Je umístěna přímo na čipu GPU.

•

CUDA – pomocí C for CUDA

Obě tyto technologie jsou určené pro podobný typ výpočtů – vektorové zpracování obrazu a dalších signálů, princip jejich funkce je však mírně odlišný. SSE jsou rozšířené instrukce procesoru pracující na principu SIMD – Single Instruction Multiple Data – Pomocí jedné instrukce je zpracování několik bloků dat. CUDA je často označována jako MIMD, respektive MSIMD (Multiple Instruction Multiple Data / Multiple SIMD).

3.3.2. L OCAL MEMORY Je to lokální paměť vlákna, která je uložená v paměti VRAM1

Překladač sám rozhodne, zda data uloží do registrů, nebo do lokální paměti, podle velikosti objektů. Která paměť byla vybrána, lze zjistit zpětnou analýzou výsledného kódu.

Pro zajímavost je rychlost porovnána také s běžnými programovacími jazyky.

3.3.3. S HARED MEMORY

•

C++

Shared memory je paměť sdílená mezi všemi vlákny uvnitř jednoho bloku.

•

C#

__shared__ float pamet[32];

Program byl přeložen jako 32bitový v prostředí .NET 3.5 ve Visual Studiu 2008 Součásti programu: • Uživatelské prostředí v C# • Výpočty v C# • DLL knihovna pro výpočty v C++ • DLL knihovna pro výpočty pomocí CUDA přeložená kompilátorem nvcc (32bit)

3.3.4. G LOBAL MEMORY Je to největší paměť, zároveň také nejpomalejší, protože není cachovaná. Je uložená v paměti VRAM. Jsou zde uloženy parametry volání kernelu a lze ji vytvořit i staticky. __device__ float pamet[32];

3.3.5. C ONSTANT MEMORY

4.2. GENEROVÁNÍ SIGNÁLU

Jde o způsob přístupu ke globální paměti. Je tedy velmi velká, ale na rozdíl od ní je cachovaná. To může při správně napsaném programu vést až k rychlosti odpovídající registrům

První signál je uložen v poli bezznaménkových jednobajtových čísel (C# byte[], C++ unsigned char *). Je vygenerován potřebný počet náhodných čísel.

3.3.6. T EXTURE MEMORY

Druhý signál je zkopírován z prvního (autokorelace), ovšem do pole, které je třikrát větší (je vložen na střed pole). Zjednoduší se tím posouvání těchto dvou signálů.

Je podobně jako constant memory pouze způsob čtení z globální paměti. Je také cachovaná. Navíc však umožňuje čtení pomocí normalizovaných souřadnic – 0 – začátek, 0.5 – polovina, 1.0 konec pole. Kromě toho umožňuje zapojit i obvody mikroprocesoru, které provádí interpolaci dat. Můžeme tak pole indexovat nejen celými čísly bez ztráty výkonu.

Výsledek je ukládána do pole dat typu int o dvojnásobné velikosti oproti délce signálu. Takovýto způsob rozložení dat v paměti odpovídá vyhledávání známého tvaru v přijatém signálu.

Signál 1

pole[10] = 5,7; pole[11] = 6,0; čtení z pole[10,5] vrátí hodnotu 5,85;

Signál 2

4. VÝPOČET KORELACE NA PC V této části bude popsán program, který byl použit pro porovnání rychlostí několika technologií na příkladu výpočtu korelace.

Korelogram Obr. č. 3: Rozložení signálu v paměti

1

VRAM = RAM grafické karty. 13-4

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

Namísto cyklů for je zde použito smyček while, které podávají výrazně lepší výkonnostní výsledky.

4.3. VÝPOČET KORELACE Z nákresů je zřejmé, že druhý signál je nutné posouvat o dvojnásobnou délku tohoto signálu. -

Při nulovém posunutí bude výsledek korelace 0 – První signál navazuje na druhý

-

Při posunutí odpovídající délce druhého signálu bude korelace nejvyšší

-

Při posunutí o dvě délky bude výsledek 0 – oba signály na sebe budou opět navazovat

4.3.2. C++ S POUŽITÍM DVOU VLÁKEN Výpočet je rozdělen do dvou současně běžících vláken, z nichž každé spočítá polovinu hodnot výsledného pole. Program poté počká na ukončení obou vláken a z výsledků vybere celková maximum. Kód je velmi podobný jako v předcházejícím příkladě.

4.3.3. C++ S POUŽITÍM SSE

C#

Nejvyššího výkonu na CPU je možné dosáhnout pouze použitím speciálních instrukcí. V tomto případě je vhodně použít rozšíření SSE, které umožňuje počítat se 4D vektory datových typů float.

Výpočet je prováděn funkcí Correlate_CSharp() v souboru Form1.cs. for (int T = 0; T < 2*Sig1.Length; T++) { for (int i=0; i < Sig1.Length; i++) { cor[T] += Sig1[i]*Sig2[i+T]; }

Kód je téměř stejný jako při použití běžného C++. Pouze výpočet vys+=(*(S1++))*(*(S2++)) je nahrazen kódem v asembleru. _asm { //S1++ movups add //S2++ movups add

if (cor[T] > cor[cor_max_index]) cor_max_index = T; } Vnější smyčka provádí časové posouvání druhého signálu, zatímco vnitřní smyčka provádí násobení jednotlivých vzorků prvního a druhého posunutého signálu. Podmínka na konci funkce hledá nejvyšší výsledek – maximum korelační funkce.

xmm0, [eax] eax, 4*4 xmm1, [ebx] ebx, 4*4

//* mulps

xmm0, xmm1

addps

xmm7, xmm0

//+=

Výpočet byl také proveden v jazyce C# pomocí rozšíření Parallel extensions a pomocí thread poolu.

//while(i2++ >= 0) sub }

4.3.1. C++ Tento jazyk je nejčastěji používán pro náročné výpočty, a byl proto v tomto měření zvolen jako referenční.

i2, 1*4

SSE umožňuje zpracovat současně až čtyři hodnoty typu float. Proto je možné pohybovat se v poli ne po jednom prvku, ale po čtyřech (instrukce add a sub)

Kód je uložen v souboru main.cpp ve funkci Correl_C while (--i1 >= 0) { int i2 = S1Len; float vys = 0.0f; float *S1 = Sig1; float *S2 = Sig2 + t; while(--i2 >= 0) { vys += (*(S1++))*(*(S2++)); } *(res++)=vys; if (vys > max_vys) { max_vys = vys; cor_max_index = t; } t++; } return cor_max_index;

4.3.4. C++ S POUŽITÍM DVOU VLÁKEN A SSE Kód je kombinací předchozího příkladu a verze C++ se dvěma vlákny.

4.3.5. C FOR CUDA Kód tohoto výpočtu je uložen v samostatné knihovně pojmenované GPUCorLib.dll. Během testů se ukázalo, že první spuštění výpočtu na GPU vytváří režii asi 150ms1. Proto je před prvním použitím z hlavního programu zavolána funkce GPUcorlib_load, které inicializuje CUDA.

1

13-5

Na testovací sestavě

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

Kód běžící na CPU

Kód běžící na GPU

Correl_GPU()

GPU_Kernel()

unsigned char *dSig1; unsigned char *dSig2; unsigned int *dOut;

unsigned int tid = (blockIdx.y*gridDim.x+blockIdx.x) * blockDim.x + threadIdx.x;

cudaMalloc((void**)&dSig1, S1Len*sizeof(char)); cudaMalloc((void**)&dSig2, S1Len*3*sizeof(char)); cudaMalloc((void**)&dOut, S1Len*2*sizeof(int));

if (tid < 2*S1Len) { int res = 0; for (int i = 0; i <S1Len; i++) { res += tex1Dfetch(tSig1,i) * tex1Dfetch(tSig2, i + tid); } Res[tid] = res; }

cudaMemcpy( dSig1, Sig1, S1Len*sizeof(char), cudaMemcpyHostToDevice) ; cudaMemcpy( dSig2, Sig2, S1Len*3*sizeof(char), cudaMemcpyHostToDevice);

Proměnná tid je index vlákna a odpovídá časovému posunutí, které je potřeba vypočítat. Protože vláken běží více, než je potřebný počet posunutí (viz funkce ceil výše), je nutné ještě před výpočtem ověřit limit

cudaBindTexture(0, tSig2, dSig2, S1Len*3*sizeof(char)); cudaBindTexture(0, tSig1, dSig1, S1Len*sizeof(char));

if (tid < 2*S1Len) Pro čtení z paměti je použita texturovací funkce tex1Dfetch s parametrem textury a souřadnice (nenormalizovaná).

int x = ceil(sqrt((2.0*S1Len)/512)); dim3 grid( x, x);

Pro ukládání vstupních dat jsou, jak již bylo řečeno, použity textury. To by mohlo být výhodné například k počítání lineárních přechodů mezi jednotlivými vzorky (počítá HW bez snížení rychlosti). V tomto případě je textura použita proto, že způsobuje výrazné zvýšení rychlosti. Projevuje se zde totiž cachování paměti, které je možné naplno využít díky tomu, že čtení z pamětí probíhá lineárně (for (int i, …, i++)). Rozdíl mezi globální a texturovou pamětí je asi v šesti řádkách kódu navíc. Rozdíl výkonu je však v tomto případě cca třicetinásobný.

testKernel<<< grid, 512 >>>(dOut, S1Len); cudaMemcpy( Res, dOut, S1Len*2*sizeof(int), cudaMemcpyDeviceToHost); cudaFree(dSig1); cudaFree(dSig2); cudaFree(dOut);

V tomto kódu nemá smysl nahrazovat smyčky for cyklem while, neboť rozdíl ve výkonu je nulový.

int max_index = 0; for (int i = 0; i < 2*S1Len; i++) { if (Res[max_index] < Res[i]) max_index = i; } return max_index;

4.4. NAMĚŘENÉ HODNOTY 4.4.1. K ONFIGURACE PŘEKLADAČŮ A TESTOVACÍ SESTAVA

Důležité parametry překladače C#:

Téměř celá část kódu je nutná proto, že GPU nemůže přistupovat do RAM. Jedná se o alokaci paměti ve VRAM a zkopírovaní dat a poté přesunutí výsledku zpět do RAM. ceil(sqrt((2.0*S1Len)/512));

/noconfig /unsafe+ /platform:x86 /debug:pdbonly /filealign:512 /optimize+ /target:winexe

spočítá velikost gridu (je použit 2D)- Velikost bloku je jednorozměrná a to 512.

Důležité parametry překladače C++ /O2 /Oi /GL /D "NDEBUG /D "UNICODE" /MD /Gy /Zi -

Důležité parametry překladače C for CUDA

-Xcompiler,/Od,/Zi,/RTC1 13-6

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

Měření probíhalo na následující sestavě:

4.4.3. V ÝSLEDKY

Intel Core 2 Duo E6550 2x2,33GHz 4GB RAM nVidia GeForce 8600GT 256MB/128bit Windows Vista 64bit ovladače GK verze 181.20 CUDA Toolki 2.2 32bit .NET Parallel Extension CTP june 08

Výpočet korelace je velmi náročný především na čtení paměti. Z toho důvodu se projevila texturová cache u GPU jako největší zrychlení. Z klasických jazyků se umístil nejlépe přístup využívající jazyk C++ a instrukční sadu SSE. Jako referenční však byl zvolen jazyk C++ s použitím jednoho vlákna. Nejpomalejší byl jazyk C#, proto, že při každém čtení ověřuje přístup do paměti. Rozšíření Parallel Extensions, které by mělo být součástí platformy .NET 4.0 podávalo velmi nestabilní výsledky a i přes velké zatížení CPU nedokázalo překonat jednovláknové C++.

4.4.2. Z PŮSOB MĚŘENÍ

Doba výpočtu [s]

Hlavní program a uživatelské rozhraní je napsáno v jazyku C# na platformě .NET 3.5. Výpočty v C++ jsou v samostatné knihovně corlib.dll a výpočty pomocí CUDA v knihovně GPUcorlib.dll. Obě knihovny byly před měřením spuštěny s nulovými daty. Zajistilo se tak jejich načtení a inicializace CUDA architektury. Každé měření bylo spuštěno desetkrát, bezprostředně po sobě. V tabulce je uvedena průměrná hodnota. Čas byl měřen pomocí proměnné Environment.TickCount. Její přesnost není nijak vysoká (+/- 10ms), ale pro tyto účely (měření v řádech vteřin až minut) je dostačující.

C++ C++ Threads C++ SSE C++ SSE Threads C# C# Parallel Ext. C# Thread Pool C for CUDA

5 0,22 0,11 0,06

10 0,81 0,42 0,20

15 1,83 0,97 0,64

20 3,24 1,70 0,81

0,05 0,53 0,44 0,30 0,02

0,20 2,04 1,56 1,19 0,05

0,66 4,60 2,54 2,56 0,10

0,81 8,14 5,73 4,42 0,17

Architektura CUDA bez použití textur podávala pouze o něco málo lepší výsledky než C++ s vlákny (cca 1,1x rychlejší). Po nahrazení klasického přístupu do paměti se rychlost zvýšila na více než třicetinásobek. Při korelaci pole o 5000 prvcích se ještě projevuje režie – tj. kopírování dat a programu do GPU. Od většího množství vzorků je však výkon již stabilní. Protože veškeré výpočty běží na grafické kartě bez zásahu hlavního procesoru, je možné CPU použít pro jakoukoliv jinou činnost, například příprava dalších dat pro výpočty, odesílání/zobrazení výsledků apod.

Počet vzorků v tisících 25 30 35 40 5,13 7,29 9,88 12,90 2,75 3,90 5,29 7,71 1,26 1,87 2,48 3,28 1,23 12,78 6,94 6,83 0,27

1,83 18,44 13,28 11,46 0,39

2,45 25,01 17,53 14,12 0,53

3,37 33,26 24,04 19,69 0,70

45 16,52 9,34 4,13

50 20,09 10,98 5,12

60 28,86 16,24 7,32

4,31 42,53 30,48 23,90 0,86

5,10 52,81 38,56 30,33 1,07

7,22 73,99 47,11 39,95 1,53

Rychlost výpočtu

Tab. č. 1: Doba výpočtu

C++ C++ Threads C++ SSE C++ SSE Threads C# C# Parallel Ext. C# Thread Pool C for CUDA

5 1,0 2,0 3,5 4,7 0,4 0,7 0,7 11,7

10 1,0 1,9 4,0 4,0 0,4 0,5 0,7 17,3

15 1,0 1,9 2,9 2,8 0,4 0,7 0,7 18,3

20 1,0 1,9 4,0 4,0 0,4 0,6 0,7 18,7

Počet vzorků v tisících 25 30 35 1,0 1,0 1,0 1,9 1,9 1,9 4,1 3,9 4,0 4,2 4,0 4,0 0,4 0,4 0,4 0,7 0,5 0,6 0,8 0,6 0,7 18,8 18,8 18,6

Tab. č. 2: Rychlost výpočtu vzhledem k C++ 13-7

40 1,0 1,7 3,9 3,8 0,4 0,5 0,7 18,5

45 1,0 1,8 4,0 3,8 0,4 0,5 0,7 19,1

50 1,0 1,8 3,9 3,9 0,4 0,5 0,7 18,8

60 1,0 1,8 3,9 4,0 0,4 0,6 0,7 18,8

2010/13 – 4. 3. 2010

VOL.12, NO.2, APRIL 2010

5. ZÁVĚR

LITERATURA

Zpracování dat na grafických kartách je velmi mladá technologie, avšak už dnes nabízí velmi vysoký výkon i při použití algoritmů stejných jako pro sériové zpracování dat. Přestože je programování na první pohled velmi jednoduché, pro skutečně maximální výkon je nutné brát v potaz mnoho detailů, jako zarovnání paměti, možnost využití cache, mnohonásobný přístup do paměti, použití registrů, rozdělení vláken do bloků/gridů a podobně. Dále je také nutné si uvědomit, že výpočet neprobíhá stejně, jako na CPU - tj. sériově, případně jako Task-Parallel, ale paralelně – Data-Parallel. Tomu je potřeba přizpůsobit i algoritmy a velmi často také data pro výpočty připravit již pomocí CPU.

[1] Intel Has Double Vision: First Multi-Core Silicon Production Begins. Intel Press Room. [Online] 7. 2 2005. [Citace: 3. 9 2009.] . [2] Questions & Misconceptions About NVIDIA GPUs and Direct3D. NVIDIA Developer Web Site. [Online] 10. 5 2004. [Citace: 3. 9 2009.] . [3] GPGPU.org. http://gpgpu.org/about. [Online] [Citace: 3. Září 2009.] .

Tento článek byl zaměřen1 na technologii CUDA. I ostatní výrobci však mají v tomto oboru své zástupce – např. ATI Stream. V době psaní tohoto článku byla skupinou Khronos vydána specifikace OpenCl[6]. Toto aplikační prostředí je určeno pro programování paralelních výpočtů. Umožňuje využít technologii nVidia CUDA i ATI Stream. Není však zaměřeno pouze na grafické akcelerátory, ale na všechny procesory umožňující paralelní výpočty(od mobilních telefonů až po servery zapojené do clusterů). Důležité je také to, že pro vytvoření OpenCl programu není nutné použít žádný speciální překladač.

[4] NVIDIA Corporation. NVIDIA CUDA™ Programming Guide 2.1. Santa Clara : autor neznámý, 2008. [5] Prof. Ing. Zdeněk Smékal, CSc. Číslicové zpracování signálů. Brno : Fakulta elektrotechniky a komunikačních technologií Vysoké učení technické v Brně, 2005. [6] Khronos Group. OpenCl. [Online] Únor 2009. [Citace: 28. Březen 2009.] .

Z výsledků je patrné, že všechny technologie/programovací jazyky si udržují konstantní rychlost, nezávisle na počtu vzorků signálu. Pouze výsledky knihovny parallel extensions kolísají. To však může být způsobeno tím, že se zatím stále nejedná o finální verzi produktu. Použitý program vytvořený ve Visual Studio 2008 Professional je možné stáhnout ze stránky http://www.moron.asp2.cz/GPU/AutoCorel.zip.

1

Z důvodu dostupného HW 13-8