Uživatel:Ladysmar

Z HPM wiki
Přejít na: navigace, hledání
Obr. 1.0: CUDA

Obsah

Historie

Technologii představila společnost NVIDIA v roce 2006. V roce 2007 byla uvolněno SDK ve verzi 1.0 pro karty NVIDIA Tesla založené na architektuře G80. Ještě v prosinci téhož roku vyšla verze CUDA SDK 1.1, která přidala podporu pro karty série GeForce verze 8. Se správným ovladačem grafické karty přibyla podpora pro překrývání paměťových přenosů výpočtem a podpora pro více GPU akcelerátorů. V roce 2008 bylo vydáno současně s architekturou G200 SDK 2.0. Postupně s verzemi SDK 2.0 - 2.3 přibývala podpora pro emulovaný výpočet v double precision a podpora pro C++ šablony v rámci kernelu. V roce 2010 je v souvislosti s architekturou Fermi vydáno SDK 3.0, kde je již nativní podpora pro double precision výpočty, podpora pro ukazatele na funkce a podpora rekurze. Vylepšeny jsou také profilovací nástroje a debuggery pro CUDA / OpenCL. Nejnovější stabilní verze CUDA SDK 4.0 byla vydána v květnu roku 2011. Největší změna je zde unifikace paměťových prostorů a masivní podpora MultiGPU.

nVidia CUDA

Uplynulo již mnoho času od dob, kdy začaly GPU (graphics compute unit) výkonnostně předhánět klasická CPU (computer processor unit), jak demonstruje první obrázek. Proto není divu, že vznikly iniciativy, aby GPU nebyly využívány jen pro zpracování 2D nebo 3D grafiky, ale pro celou škálu jiných, rovněž výkonnostně náročných aplikací a jejich potenciál, kterým je především kvalitní podpora pro paralelní výpočty na multiprocesorech, byl efektivně využit.


Úvod do technologie CUDA

Z těchto iniciativ pak vznikly konkrétní technologie, jako Compute Unified Device Architecture (dále jen CUDA) společnosti nVidia nebo ATI Stream společnosti AMD/ATI či standard Open Computing Language (dále jen OpenCL) pod záštitou konsorcia Khronos ,tvořená společnostmi jako AMD, Intel, nVidia, IBM, a dalších 100 společností, která rovněž udržuje dobře známý standard OpenGL.

Jelikož není v českých končinách ani jedna z výše popisovaných technologií podrobněji popsána (snad jen kromě menších zpráviček a akademické půdy),rozhodl jsem se napsat seriál o základním popisu technologie CUDA.


Srovnání výkonu nVidia GPU s Intel CPU

Obr. 1.1: Srovnání výkonu nVidia GPU s Intel CPU

Základní popis a architektura CUDA

A nyní si povíme něco o základní architektuře CUDA. Nyní si na obrázku 1.2 popíšeme různé vrsty architektury CUDA a začneme jazyky, kde je zřejmé, že CUDA podporuje různé jazyky a API, jako Fortran, OpenCL, C a nebo DirectX Computer. ovšem pro nás je důležité a nejvíce známí jazyk C. O úroveň výše je vidět, že CUDA podporuje všechny dostupné operační systémy, OS X, Windows, Linux. Rovněž si ukážeme i všechny odvětví kam se CUDA hodí, jako je například věda, medicína, zpracování videa,film, offline rendring apod. kde je potřeba vynikající grafika.

Obr. 1.2:Architektura CUDA
CPU vs. GPU


A nyní si ukážeme jak vypadá CUDA zevnitř (z pohledu GPU)? To už nám prozrazuje obr. 1.3, kde můžeme vidět, že celé zařízení je rozděleno na tyto části: multiprocesory (například GeForce 8800 GTS jich má 12) – každý multiprocesor má 8 procesorů instrukční jednotka procesory – 32-bitový procesor architektury SIMT (Single Instruction Multi-Thread vycházející z architektury SIMD) a jak bylo uvedeno výše, 8 procesorů tvoří jeden multiprocesor, registry – 32 bitové, 8192 registrů na G80, sdílená paměť – kešovaná, velikost 16KB na G80, paměť pro konstanty – kešovaná, paměť pro textury – kešovaná, globální paměť – paměť určená pro kopírování dat z RAM do GPU a naopak.

Obr. 1.3:Architektura CUDA z pohledu GPU

Jak je vidět z dalšího obrázku (1.4),CUDA aplikace může přistupovat k jakékoliv části API CUDA. Ať už se jedná o knihovny CUBLAS (Basic Linear Algebra Subprograms) nebo CUFFT (Fast Fourier Transform), tak i runtime API nebo driver API.

Obr. 1.4:Přístup programu k API CUDA


Struktura multiprocesoru

Obecně se multiprocesor skládá z několika (dnes až ze 32) stream procesorů, pole registrů, sdílené paměti, několika load/store jednotek a Special Function Unit - jednotky pro výpočet složitějších funkcí jako sin,cos,ln.


Schéma streaming multiprocesoru architektury Fermi



Sdílení paměti

O sdílené paměti víme, že je velmi rychlá (v oficiální dokumentaci se uvádí rychlost srovnatelná s registry), jak ji alokovat pro naše programy a že je dobrým pomocníkem při sdruženém přístupu do globální paměti (pokud není dodržena velikost paměťového elementu na 4, 8 nebo 16 bytech). Co je důležité vědět o sdílené paměti? Se sdílenou pamětí je provázána existence bank a konfliktů bank.

Banky jsou malé paměťové elementy (přístup do banky zabere 2 cykly) o velikosti 32bitů (4 byty).Nedojde-li ke konfliktu bank, lze k nim přistupovat současně prostřednictvím half-warpu. V rámci celého warpu se jedná o 2 “transakce”, kde jedna je pro dolní a druhá pro horní half-warp. Konflikt bank spočívá v tom, že dvě paměťové adresy v rámci half-warpu odkazují do stejné banky. V takovém případě se přístup do bank serializuje a není již současný. Konfliktu bank se říká n-pásmové bankovní konflikty (n-way bank conflicts), kde n označuje počet paměťových adres přistupujících do jedné banky.

Existuje jedna vyjímka, kdy nedojde ke konfliktu bank, a to když přistupuje všech 16 vláken half-warpu ke stejné bance.

Nyní si ukážeme grafické znázornění přístupu do bank s/bez konfliktů.

Obr. 1.5:Přístup do bank sdílené paměti bez konfliktu bank
Obr. 1.6:Přístup do bank sdílené paměti s konfliktem bank


Programovací model

Uspořádání vláken a bloků

CUDA aplikace je složena z částí, které běží buď na hostu (CPU) nebo na CUDA zařízení (GPU). Části aplikace běžící na zařízení jsou spouštěny hostem zavoláním kernelu, což je funkce, která je prováděna každým spuštěným vláknem (thread).

Blok (thread block
Vlákna jsou organizována do 1D, 2D nebo 3D bloků, kde vlákna ve stejném bloku mohou sdílet data a lze synchronizovat jejich běh. Počet vláken na jeden blok je závislý na výpočetních možnostech zařízení. Každé vlákno je v rámci bloku identifikováno unikátním indexem přístupným ve spuštěném kernelu přes zabudovanou proměnou threadIdx.
Mřížka (grid
Bloky jsou organizovány do 1D, 2D nebo 3D mřížky. Blok lze v rámci mřížky identifikovat unikátním indexem přístupným ve spuštěném kernelu přes zabudovanou proměnou blockIdx. Každý blok vláken musí být schopen pracovat nezávisle na ostatních, aby byla umožněna škálovatelnost systému (na GPU s více jádry půjde spustit více bloků paralelně oproti GPU s méně jádry kde bloky poběží v sérii).
Warp 
Balík vláken zpracovávaných v jednom okamžiku se nazývá warp. Jeho velikost je závislá na počtu výpočetních jednotek.

Počet a organizace spuštěných vláken v jednom bloku a počet a organizace bloků v mřížce se určuje při volání kernelu.

Typický průběh GPGPU výpočtu

  1. Vyhrazení paměti na GPU
  2. Přesun dat z hlavní paměti RAM do paměti grafického akcelerátoru
  3. Spuštění výpočtu na grafické kartě
  4. Přesun výsledků z paměti grafické karty do hlavní RAM paměti

Ukázka kódu v CUDA C

Převzato z wikipedia.org

// Kód pro GPU__global__ void VecAdd(float* A, float* B, float* C, int N){

   int i = blockDim.x * blockIdx.x + threadIdx.x;
   if (i < N)
       C[i] = A[i] + B[i];

} // Kód pro CPU int main() {

   int N = ...;
   size_t size = N * sizeof(float);

// Alokace vstupních vektorů h_A and h_B v hlavní paměti

   float* h_A = (float*)malloc(size);
   float* h_B = (float*)malloc(size);

// Inicializace vstupních vektorů

  ...

// Alokace paměti na zařízení

   float* d_A;
   cudaMalloc(&d_A, size);
   float* d_B;
   cudaMalloc(&d_B, size);
   float* d_C;
   cudaMalloc(&d_C, size);

// Přesun vektorů z hlavní paměti do paměti zařízení

   cudaMemcpy(d_A, h_A, size, cudaMemcpyHostToDevice);
   cudaMemcpy(d_B, h_B, size, cudaMemcpyHostToDevice);

// Volání kernelu

   int threadsPerBlock = 256;
   int blocksPerGrid = (N + threadsPerBlock – 1) / threadsPerBlock;
   VecAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, N);

// Přesun výsledků z paměti zařízení do hlavní paměti

   cudaMemcpy(h_C, d_C, size, cudaMemcpyDeviceToHost);

// Uvolnění paměti na zařízení

   cudaFree(d_A);
   cudaFree(d_B);
   cudaFree(d_C);

// Uvolnění hlavní paměti ... </source>

CUDA Visual Profiler

CUDA Visual Profiler je nástrojem sloužícím k přehledu všech možných informácí o CUDA programu. Zde je výčet hlavních informací, které můžeme díky Visual Profileru sledovat: sdružené čtení z globální paměti (gld coalesced)sdružený zápis z globální paměti (gst coalesced)nesdružené čtení z globální paměti (gld uncoalesced)nesdružený zápis z globální paměti (gst uncoalesced)x složka gridu (grid size X)y složka gridu (grid size Y)

A teď pár informací pro koncové uživatele

První věc, která vás na prezentaci nVidie vyskočí je podpora rozhraní CUDA.Mobilní GeForce od řady 8000M dále nyní podporují kompletně celý programovací jazyk CUDA. Co to znamená pro vás jako pro koncové uživatele? Rychlejší enkódování videa, možnost aplikace různých filtrů při sledování videa, které spočítá grafická karta... Vlastně nic zásadního. Tedy zásadní to bude pro programátory, kteří budou moci pro programování více využívat i běžných notebooků a ne jen ty s mobilními nVidia Quadro, ale to se nás, běžných uživatelů, příliš netýká.

Badaboom

Obr. 1.7:Badaboom


Vlastníte kromě notebooku s nVidia GeForce kartou ještě iPhone a rádi cestujete? Pak CUDA více potěší alespoň vás, protože díky speciálnímu programu můžete videa konvertovat do iPhonu přes grafickou kartu, což dle tvrzení nVidie trvá podstatně méně času.

Zajímavost na konec

super počítač

pojmy v tomto textu:

FLOPS je zkratka pro počet operací v plovoucí řádové čárce za sekundu (FLoating-point OPerations per Second)

Zde najdete zajímavé video o tomto super počítači: http://pcworld.cz/hardware/superpocitac-nvidia-tesla-se-predstavuje-svetu-3338

Tesla super počítač

Superpočítač s označením Tesla poprvé veřejně představil v Londýně společnost nVidia. Jedná se o doopravdy neobyčejný počítač, jelikož jeho počáteční výkon nezajišťují klasické procesory CPU, ale GPU - tedy extrémně výkonné čipy, navržené hlavně pro podporu grafiky u 3D počítačových her. Mezi nesporné výhody architektury patří paralelní zpracování úkolů. V klaciských počítačích se používá maximálně několik procesorů, v případě Tesly, jde až o 960 procesorů s výkonem až 4 TFLOPs a tím pádem je výkon Tesly až 250x vyšší než výkon osobního počítače.

Výhodou Nvidia CUDA systému je fakt, že nikdo nenutí vědce, aby se učili speciální postupy nebo nové programovací jazyky. I při využití Tesly tak mohou jednoduše využívat rozšířený jazyk C. Do budoucna by se měla objevit i podpora pro programovací jazyky C++ a Fortran.

Díky CUDA GPU se výpočty lékařských modelů na z týdne stáhli na 51 minut a to díky pouhými 8 grafickými kartami GPU. Studenti udělali pokus tento počítač vs. CPU složený z 512ti operonů a i když zkrátili čas, tak ani grafikou a ani časovou osou se nedostali na stejnou rychlost jako GPU nVidia CUDA.

jako další zajímavý příklad můžu uvést příklad výpočtu průtoku vody turbínou na počítači s 2,5 GHz, kde výpočet trval celkem 12 hodin a při výpočtu na superpočítači Tesla s 4 GPU pouze 9 minut.

Nejvýkonnější superpočítač

Tento cluster je složený z 32 Tesla super počítačů s neuvěřitelným výkonem 128 TFLOPs.

Je vidět, že to nemusejí být obrovské sálové mašiny, nebo stovky propojených počítačů do výpočetních clusterů, které mají obrovskou spotřebu energie. Společnost Nvidia dnes ukázala jiný přístup, zdá se, že mnohem elegantnější a efektivnější i ekonomičtější.

Tesla se proto uplatní především v odvětvích, které jsem ve svém článku zmínil, při sekvenčních výpočtech, kdy se následně kalkuluje s předchozími výsledky výpočtů, je urychlení mnohem menší. Nicméně, i přesto je přínos Tesly markantní.

Video

grafické zpracování pomocí nVidia Cuda

NVIDIA CUDA realtime particles demo

https://www.youtube.com/embed/RqduA7myZok

Sony Vegas 10 + nVIDIA CUDA

http://www.youtube.com/embed/cYxrOnN5_5o

Závěrem

Vědci nyní budou moci dělat extrémně náročné výpočty, simulace a pokusy rovnou ve své laboratoři či dokonce doma. Za tento luxus mohou poděkovat nejen vývojářům z Nvidie, ale paradoxně i počítačovým hráčům, neboť bouřlivý vývoj GPU procesorů financovali z velké části právě oni při honbě za nejlepší grafikou ve svých oblíbených hrách.

Reference

CUDA - wikipedia.org

zdroj použití videa

Popis CUDA včetně příkladů pro programátory a nastavení převážně Linux

Super počítač nVidia Tesla

Osobní nástroje
Jmenné prostory
Varianty
Akce
Navigace
Nástroje