Podívejte se, co je „CUDA“ v jiných slovnících. Počítačový zdroj U SM Programy využívající cuda

A je navržen tak, aby překládal hostitelský kód (hlavní, řídicí kód) a kód zařízení (hardwarový kód) (soubory s příponou .cu) do objektových souborů vhodných pro proces sestavení finálního programu nebo knihovny v jakémkoli programovacím prostředí, např. v NetBeans.

Architektura CUDA využívá model mřížkové paměti, modelování vláken clusteru a instrukce SIMD. Použitelné nejen pro vysoce výkonné grafické výpočty, ale také pro různé vědecké výpočty pomocí grafických karet nVidia. Vědci a výzkumníci široce používají CUDA v různých oblastech, včetně astrofyziky, výpočetní biologie a chemie, modelování dynamiky tekutin, elektromagnetických interakcí, počítačové tomografie, seismické analýzy a dalších. CUDA má schopnost připojit se k aplikacím pomocí OpenGL a Direct3D. CUDA je multiplatformní software pro operační systémy jako Linux, Mac OS X a Windows.

22. března 2010 nVidia vydala CUDA Toolkit 3.0, který obsahoval podporu pro OpenCL.

Zařízení

Platforma CUDA se poprvé objevila na trhu s vydáním čipu NVIDIA G80 osmé generace a stala se přítomnou ve všech následujících řadách grafických čipů, které se používají v rodinách akcelerátorů GeForce, Quadro a NVidia Tesla.

První řada hardwaru s podporou CUDA SDK, G8x, měla 32bitový vektorový procesor s jednoduchou přesností využívající CUDA SDK jako API (CUDA podporuje typ C double, ale jeho přesnost byla nyní snížena na 32bitové plovoucí bod). Pozdější procesory GT200 mají podporu pro 64bitovou přesnost (pouze SFU), ale výkon je výrazně horší než u 32bitové přesnosti (kvůli skutečnosti, že na stream multiprocesoru jsou pouze dvě SFU, zatímco skalárních procesorů je osm). GPU organizuje hardwarové multithreading, což vám umožňuje využívat všechny zdroje GPU. Otevírá se tak možnost přenést funkce fyzického akcelerátoru do grafického akcelerátoru (příkladem implementace je nVidia PhysX). Otevírá také široké možnosti využití grafický hardware počítač pro provádění složitých negrafických výpočtů: například ve výpočetní biologii a dalších vědních oborech.

Výhody

Ve srovnání s tradičním přístupem k organizování výpočtů obecný účel Díky možnostem grafických rozhraní API má architektura CUDA v této oblasti následující výhody:

Omezení

  • Všechny funkce spustitelné na zařízení nepodporují rekurzi (CUDA Toolkit 3.1 podporuje ukazatele a rekurzi) a mají některá další omezení

Podporované GPU a grafické akcelerátory

Seznam zařízení od výrobce zařízení Nvidia s deklarovanou plnou podporou technologie CUDA je uveden na oficiálních stránkách Nvidie: CUDA-Enabled GPU Products (anglicky).

Ve skutečnosti následující periferie aktuálně podporují technologii CUDA na trhu s PC hardwarem:

Verze specifikace GPU Video karty
1.0 G80, G92, G92b, G94, G94b GeForce 8800GTX/Ultra, 9400GT, 9600GT, 9800GT, Tesla C/D/S870, FX4/5600, 360M, GT 420
1.1 G86, G84, G98, G96, G96b, G94, G94b, G92, G92b GeForce 8400GS/GT, 8600GT/GTS, 8800GT/GTS, 9600 GSO, 9800GTX/GX2, GTS 250, GT 120/30/40, FX 4/570, 3/580, 3700x1x8,3 /370M, 3/5/770M, 16/17/27/28/36/37/3800M, NVS420/50
1.2 GT218, GT216, GT215 GeForce 210, GT 220/40, FX380 LP, 1800M, 370/380M, NVS 2/3100M
1.3 GT200, GT200b GeForce GTX 260, GTX 275, GTX 280, GTX 285, GTX 295, Tesla C/M1060, S1070, Quadro CX, FX 3/4/5800
2.0 GF100, GF110 GeForce (GF100) GTX 465, GTX 470, GTX 480, Tesla C2050, C2070, S/M2050/70, Quadro Plex 7000, Quadro 4000, 5000, 6000, GTX 480, GTX580, GTX54TX 560 0
2.1 GF104, GF114, GF116, GF108, GF106 GeForce 610M, GT 430, GT 440, GTS 450, GTX 460, GTX 550 Ti, GTX 560, GTX 560 Ti, 500M, Quadro 600, 2000
3.0 GK104, GK106, GK107 GEFORCE GTX 690, GTX 680, GTX 670, GTX 660 Ti, GTX 660, GTX 650 Ti, GTX 650, GT 640, GeForce GTX 680MX, GeForce GTX 680M, GTX 675, MXTX, GeForce, GTX 675MX, GeForce GTX 670MX, GTX 670MX, GTX 670MX, GTX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 670MX, GTX 70MX, GTX 670 70MX, GTX 670MX, GTX, GTX TX 660M, GeForce GT 650M, GeForce GT 645M, GeForce GT 640M
3.5 GK110
Nvidia GeForce pro stolní počítače
GeForce GTX 590
GeForce GTX 580
GeForce GTX 570
GeForce GTX 560 Ti
GeForce GTX 560
GeForce GTX 550 Ti
GeForce GTX 520
GeForce GTX 480
GeForce GTX 470
GeForce GTX 465
GeForce GTX 460
GeForce GTS 450
GeForce GTX 295
GeForce GTX 285
GeForce GTX 280
GeForce GTX 275
GeForce GTX 260
GeForce GTS 250
GeForce GT 240
GeForce GT 220
GeForce 210
GeForce GTS 150
GeForce GT 130
GeForce GT 120
GeForce G100
GeForce 9800 GX2
GeForce 9800 GTX+
GeForce 9800 GTX
GeForce 9800 GT
GeForce 9600 GSO
GeForce 9600 GT
GeForce 9500 GT
GeForce 9400 GT
GeForce 9400 mGPU
GeForce 9300 mGPU
GeForce 8800 GTS 512
GeForce 8800 GT
GeForce 8600 GTS
GeForce 8600 GT
GeForce 8500 GT
GeForce 8400GS
Nvidia GeForce pro mobilní počítače
GeForce GTX 580M
GeForce GTX 570M
GeForce GTX 560M
GeForce GT 555M
GeForce GT 540M
GeForce GT 525M
GeForce GT 520M
GeForce GTX 485M
GeForce GTX 480M
GeForce GTX 470M
GeForce GTX 460M
GeForce GT 445M
GeForce GT 435M
GeForce GT 425M
GeForce GT 420M
GeForce GT 415M
GeForce GTX 285M
GeForce GTX 280M
GeForce GTX 260M
GeForce GTS 360M
GeForce GTS 350M
GeForce GTS 160M
GeForce GTS 150M
GeForce GT 335M
GeForce GT 330M
GeForce GT 325M
GeForce GT 240M
GeForce GT 130M
GeForce G210M
GeForce G110M
GeForce G105M
GeForce 310M
GeForce 305M
GeForce 9800M GTX
GeForce 9800M GT
GeForce 9800M GTS
GeForce 9700M GTS
GeForce 9700M GT
GeForce 9650MGS
GeForce 9600M GT
GeForce 9600MGS
GeForce 9500MGS
GeForce 9500M G
GeForce 9300MGS
GeForce 9300M G
GeForce 9200MGS
GeForce 9100M G
GeForce 8800M GTS
GeForce 8700M GT
GeForce 8600M GT
GeForce 8600MGS
GeForce 8400M GT
GeForce 8400MGS
Nvidia Tesla *
Tesla C2050/C2070
Tesla M2050/M2070/M2090
Tesla S2050
Tesla S1070
Tesla M1060
Tesla C1060
Tesla C870
Tesla D870
Tesla S870
Nvidia Quadro pro stolní počítače
Quadro 6000
Quadro 5000
Quadro 4000
Quadro 2000
Quadro 600
Quadro FX 5800
Quadro FX 5600
Quadro FX 4800
Quadro FX 4700 X2
Quadro FX 4600
Quadro FX 3700
Quadro FX 1700
Quadro FX 570
Quadro FX 470
Quadro FX 380 s nízkým profilem
Quadro FX 370
Quadro FX 370 s nízkým profilem
Quadro CX
Quadro NVS 450
Quadro NVS 420
Quadro NVS 290
Quadro Plex 2100 D4
Quadro Plex 2200 D2
Quadro Plex 2100 S4
Quadro Plex 1000 Model IV
Nvidia Quadro pro mobilní počítače
Quadro 5010M
Quadro 5000M
Quadro 4000M
Quadro 3000M
Quadro 2000M
Quadro 1000M
Quadro FX 3800M
Quadro FX 3700M
Quadro FX 3600M
Quadro FX 2800M
Quadro FX 2700M
Quadro FX 1800M
Quadro FX 1700M
Quadro FX 1600M
Quadro FX 880M
Quadro FX 770M
Quadro FX 570M
Quadro FX 380M
Quadro FX 370M
Quadro FX 360M
Quadro NVS 5100M
Quadro NVS 4200M
Quadro NVS 3100M
Quadro NVS 2100M
Quadro NVS 320M
Quadro NVS 160M
Quadro NVS 150M
Quadro NVS 140M
Quadro NVS 135M
Quadro NVS 130M
  • Modely Tesla C1060, Tesla S1070, Tesla C2050/C2070, Tesla M2050/M2070, Tesla S2050 umožňují výpočty GPU s dvojnásobnou přesností.

Vlastnosti a specifikace různých verzí

Podpora funkcí (neuvedené funkce jsou
podporováno pro všechny výpočetní schopnosti)
Výpočetní schopnost (verze)
1.0 1.1 1.2 1.3 2.x

32bitová slova v globální paměti
Ne Ano

hodnoty s plovoucí desetinnou čárkou v globální paměti
Celočíselné atomické funkce fungující na
32bitová slova ve sdílené paměti
Ne Ano
atomicExch() pracující na 32bit
hodnoty s pohyblivou řádovou čárkou ve sdílené paměti
Celočíselné atomické funkce fungující na
64bitová slova v globální paměti
Funkce warp hlasování
Operace s plovoucí desetinnou čárkou s dvojitou přesností Ne Ano
Atomové funkce fungující na 64-bit
celočíselné hodnoty ve sdílené paměti
Ne Ano
Atomová adice s plovoucí desetinnou čárkou funguje
32bitová slova v globální a sdílené paměti
_hlasování()
_threadfence_system()
_syncthreads_count(),
_syncthreads_and(),
_syncthreads_or()
Funkce povrchu
3D mřížka bloku závitů
Technické specifikace Výpočetní schopnost (verze)
1.0 1.1 1.2 1.3 2.x
Maximální rozměrnost rastru závitových bloků 2 3
Maximální rozměr x, y nebo z mřížky bloků závitů 65535
Maximální rozměrnost závitového bloku 3
Maximální rozměr x nebo y bloku 512 1024
Maximální z-rozměr bloku 64
Maximální počet vláken na blok 512 1024
Velikost osnovy 32
Maximální počet rezidentních bloků na multiprocesor 8
Maximální počet rezidentních warps na multiprocesor 24 32 48
Maximální počet rezidentních vláken na víceprocesor 768 1024 1536
Počet 32bitových registrů na multiprocesor 8 tis 16 K 32 K
Maximální množství sdílené paměti na víceprocesor 16 kB 48 kB
Počet bank sdílené paměti 16 32
Množství místní paměti na vlákno 16 kB 512 kB
Konstantní velikost paměti 64 kB
Pracovní sada mezipaměti na multiprocesor pro konstantní paměť 8 kB
Mezipaměť pracovní sady na víceprocesor pro texturovou paměť V závislosti na zařízení, mezi 6 KB a 8 KB
Maximální šířka pro 1D texturu
8192 32768
Maximální šířka pro 1D texturu
reference vázaná na lineární paměť
2 27
Maximální šířka a počet vrstev
pro referenci 1D vrstvené textury
8192 x 512 16384 x 2048
Maximální šířka a výška pro 2D
odkaz na texturu vázaný na
lineární paměť nebo pole CUDA
65536 x 32768 65536 x 65535
Maximální šířka, výška a počet
vrstev pro referenci 2D vrstvené textury
8192 x 8192 x 512 16384 x 16384 x 2048
Maximální šířka, výška a hloubka
pro referenci 3D textury vázanou na lineární
paměti nebo pole CUDA
2048 x 2048 x 2048
Maximální počet textur, které
lze vázat na jádro
128
Maximální šířka pro 1D povrch
odkaz vázaný na pole CUDA
Ne
podporováno
8192
Maximální šířka a výška pro 2D
povrchová reference vázaná na pole CUDA
8192 x 8192
Maximální počet povrchů, které
lze vázat na jádro
8
Maximální počet instrukcí za
jádro
2 miliony

Příklad

CudaArray* cu_array; textura< float , 2 >tex; // Přidělení pole cudaMalloc( & cu_array, cudaCreateChannelDesc< float>(), šířka výška) ; // Zkopírujte data obrázku do pole cudaMemcpy( cudaMemcpy( cuda_array, image, width* height, cudaMemcpyHostToDevice) ; // Svázat pole s texturou cudaBindTexture( tex, cu_array) ; // Spustit kernel dim3 blockDim(16, 16, 1) ; dim3 gridDim(width / blockDim.x, height / blockDim.y, 1) ; jádro<<< gridDim, blockDim, 0 >>> (d_odata, sirka, vyska) ; cudaUnbindTexture(tex) ; __global__ void kernel(float * odata, int výška, int šířka) ( unsigned int x = blockIdx.x * blockDim.x + threadIdx.x ; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y ; float c = texfetch(tex, x, y) ; odata[ y* šířka+ x] = c; )

Import pycuda.driver jako drv import numpy drv.init() dev = drv.Device(0) ctx = dev.make_context() mod = drv.SourceModule( """ __global__ void multiply_them(float *dest, float *a, float *b) ( const int i = threadIdx.x; dest[i] = a[i] * b[i]; ) """) multiply_them = mod.get_function ("multiply_them" ) a = numpy.random .randn (400 ) .astype (numpy.float32 ) b = numpy.random .randn (400 ) .astype (numpy.float32 ) jako dest = . (a) multiply_them( drv.Out (dest) , drv.In (a) , drv.In (b) , block= (400 , 1 , 1 ) ) print dest-a*b

CUDA jako předmět na vysokých školách

Od prosince 2009 se softwarový model CUDA vyučuje na 269 univerzitách po celém světě. V Rusku probíhají školení o CUDA na Petrohradské polytechnické univerzitě, Jaroslavské státní univerzitě. P. G. Demidov, Moskva, Nižnij Novgorod, Petrohrad, Tver, Kazaň, Novosibirsk, Novosibirská státní technická univerzita, Omské a Permské státní univerzity, Mezinárodní univerzita povahy společnosti a člověka „Dubna“, Ivanovská státní energetická univerzita, Bělgorodská státní univerzita , MSTU je. Bauman, Ruská chemická technická univerzita pojmenovaná po. Mendělejev, Meziregionální superpočítačové centrum RAS, . Kromě toho bylo v prosinci 2009 oznámeno, že první ruské vědecké a vzdělávací centrum „Parallel Computing“, které se nachází ve městě Dubna, začalo fungovat, mezi jehož úkoly patří školení a konzultace o řešení složitých výpočetních problémů na GPU.

Na Ukrajině se kurzy CUDA vyučují na Kyjevském institutu systémové analýzy.

Odkazy

Oficiální zdroje

  • CUDA Zone (rusky) - oficiální stránky CUDA
  • CUDA GPU Computing (anglicky) - oficiální webová fóra věnovaná CUDA computingu

Neoficiální zdroje

Tomův hardware
  • Dmitrij Čekanov. nVidia CUDA: výpočty na grafické kartě nebo smrt CPU? . Tom's Hardware (22. června 2008). Archivováno
  • Dmitrij Čekanov. nVidia CUDA: Benchmarking GPU aplikací pro masový trh. Tom's Hardware (19. května 2009). Archivováno z originálu 4. března 2012. Získáno 19. května 2009.
iXBT.com
  • Alexey Berillo. NVIDIA CUDA - negrafické výpočty na GPU. Část 1 . iXBT.com (23. září 2008). Archivováno z originálu 4. března 2012. Získáno 20. ledna 2009.
  • Alexey Berillo. NVIDIA CUDA - negrafické výpočty na GPU. Část 2 . iXBT.com (22. října 2008). - Příklady implementace NVIDIA CUDA. Archivováno z originálu 4. března 2012. Získáno 20. ledna 2009.
Jiné zdroje
  • Boreskov Alexej Viktorovič. Základy CUDA (20. ledna 2009). Archivováno z originálu 4. března 2012. Získáno 20. ledna 2009.
  • Vladimír Frolov.Úvod do technologie CUDA. Online časopis „Počítačová grafika a multimédia“ (19. prosince 2008). Archivováno z originálu 4. března 2012. Získáno 28. října 2009.
  • Igor Oskolkov. NVIDIA CUDA je cenově dostupná vstupenka do světa velkých počítačů. Computerra (30. dubna 2009). Staženo 3. května 2009.
  • Vladimír Frolov.Úvod do technologie CUDA (1. srpna 2009). Archivováno z originálu 4. března 2012. Získáno 3. dubna 2010.
  • GPGPU.ru. Použití grafických karet pro výpočetní techniku
  • . Paralelní výpočetní centrum

Poznámky

viz také

Vraťme se do roku 2003, kdy Intel a AMD společně závodily v hledání nejvýkonnějšího procesoru. Za pouhých pár let se v důsledku tohoto závodu výrazně zvýšily takty, zejména po vydání Intel Pentium 4.

Závod se ale rychle blížil k limitu. Po vlně obrovského nárůstu taktů (mezi lety 2001 a 2003 se takt Pentia 4 zdvojnásobil z 1,5 na 3 GHz) se uživatelé museli spokojit s desetiny gigahertzů, které byli výrobci schopni vymáčknout (od roku 2003 do roku 2005 se takt zvýšil z pouhých 3 na 3,8 GHz).

I architektury optimalizované pro vysoké taktovací frekvence, jako je Prescott, se začaly potýkat s potížemi, a tentokrát nejen produkčními. Výrobci čipů prostě narazili na fyzikální zákony. Někteří analytici dokonce předpovídali, že Moorův zákon přestane platit. To se ale nestalo. Původní význam zákona je často zkreslený, ale týká se počtu tranzistorů na povrchu křemíkového jádra. Po dlouhou dobu byl nárůst počtu tranzistorů v CPU doprovázen odpovídajícím zvýšením výkonu - což vedlo ke zkreslení významu. Pak se ale situace zkomplikovala. Vývojáři architektury CPU přistoupili na zákon redukce růstu: počet tranzistorů, které bylo potřeba přidat pro požadované zvýšení výkonu, byl stále větší, což vedlo do slepé uličky.



Zatímco si výrobci CPU trhali vlasy ve snaze najít řešení svých problémů, výrobci GPU nadále pozoruhodně těžili z výhod Moorova zákona.

Proč se nedostali do stejné slepé uličky jako vývojáři architektury CPU? Důvod je velmi jednoduchý: centrální procesorové jednotky jsou navrženy pro příjem maximální výkon na proudu instrukcí, které zpracovávají různá data (jak celá čísla, tak čísla s pohyblivou řádovou čárkou), provádějí náhodný přístup do paměti atd. Až dosud se vývojáři snaží poskytovat větší paralelismus instrukcí – tedy provádět co nejvíce instrukcí paralelně. Například u Pentia se objevilo superskalární provádění, kdy za určitých podmínek bylo možné provést dvě instrukce za takt. Pentium Pro obdrželo mimo pořadí provádění instrukcí, což umožnilo optimalizovat chod výpočetních jednotek. Problém je v tom, že paralelní provádění sekvenčního proudu instrukcí má zjevná omezení, takže slepé zvyšování počtu výpočetních jednotek nepřináší žádnou výhodu, protože budou většinu času stále nečinné.

Naproti tomu obsluha GPU je poměrně jednoduchá. Skládá se z odebrání skupiny polygonů na jedné straně a generování skupiny pixelů na straně druhé. Polygony a pixely jsou na sobě nezávislé, takže je lze zpracovávat paralelně. V GPU je tedy možné velkou část krystalu alokovat do výpočetních jednotek, které se na rozdíl od CPU skutečně využijí.



Pro zvětšení klikněte na obrázek.

GPU se od CPU liší nejen tímto způsobem. Přístup k paměti v GPU je velmi svázaný - pokud je čten texel, pak po několika hodinových cyklech bude načten sousední texel; Když je zaznamenán pixel, po několika hodinových cyklech bude zaznamenán sousední pixel. Inteligentním uspořádáním paměti můžete dosáhnout výkonu blízkého teoretické propustnosti. To znamená, že GPU na rozdíl od CPU nevyžaduje velkou mezipaměť, protože jeho úlohou je urychlit operace texturování. Vše, co je potřeba, je několik kilobajtů obsahujících několik texelů používaných v bilineárních a trilineárních filtrech.



Pro zvětšení klikněte na obrázek.

Ať žije GeForce FX!

Oba světy zůstaly na dlouhou dobu odděleny. Pro kancelářské úkoly a internetové aplikace jsme používali CPU (nebo dokonce více CPU) a GPU byly dobré pouze pro urychlení vykreslování. Jedna vlastnost ale vše změnila: totiž příchod programovatelných GPU. Zpočátku se CPU neměly čeho bát. První takzvané programovatelné GPU (NV20 a R200) ​​nebyly téměř žádnou hrozbou. Počet instrukcí v programu zůstal omezen na asi 10 a pracovaly s velmi exotickými datovými typy, jako jsou 9- nebo 12bitová čísla s pevnou řádovou čárkou.



Pro zvětšení klikněte na obrázek.

Ale Moorův zákon opět ukázal svou nejlepší stránku. Zvýšení počtu tranzistorů umožnilo nejen zvýšit počet výpočetních jednotek, ale také zlepšilo jejich flexibilitu. Vzhled NV30 lze považovat za významný krok vpřed z několika důvodů. Hráčům se samozřejmě karty NV30 příliš nelíbily, ale nová GPU se začala spoléhat na dvě funkce, které byly navrženy tak, aby změnily vnímání GPU jako více než jen grafických akcelerátorů.

  • Podpora výpočtů s plovoucí desetinnou čárkou s jednoduchou přesností (i když nevyhovuje standardu IEEE754);
  • podpora více než tisíce návodů.

Máme tedy všechny podmínky, které mohou přilákat průkopnické výzkumníky, kteří neustále hledají další výpočetní výkon.

Myšlenka použití grafických akcelerátorů pro matematické výpočty není nová. První pokusy byly provedeny již v 90. letech minulého století. Samozřejmě byly velmi primitivní – většinou omezené na použití některých hardwarových funkcí, jako je rasterizace a Z-buffery, pro urychlení úloh, jako je hledání trasy nebo odvození. Voronoiovy diagramy .



Pro zvětšení klikněte na obrázek.

V roce 2003, s příchodem vyvinutých shaderů, bylo dosaženo nového baru - tentokrát provádějícího maticové výpočty. To byl rok, kdy byla celá část SIGGRAPH („Výpočty na GPU“) věnována nové oblasti IT. Tato raná iniciativa se nazývala GPGPU (General-Purpose computation on GPU). A brzkým zlomem byl vznik .

Abyste pochopili roli BrookGPU, musíte pochopit, jak se všechno stalo před jeho objevením. Jediný způsob, jak získat zdroje GPU v roce 2003, bylo použít jedno ze dvou grafických API – Direct3D nebo OpenGL. V důsledku toho se vývojáři, kteří chtěli funkce GPU pro své výpočty, museli spoléhat na dvě zmíněná API. Problém je v tom, že ne vždy byli odborníky na programování grafických karet, a to vážně komplikovalo přístup k technologii. Pokud 3D programátoři pracují se shadery, texturami a fragmenty, pak specialisté v oblasti paralelního programování spoléhají na vlákna, jádra, rozptyly atd. Proto bylo nejprve nutné nakreslit analogie mezi těmito dvěma světy.

  • Proud je proud prvků stejného typu, v GPU může být reprezentován texturou. V principu klasického programování existuje takový analog jako pole.
  • Jádro- funkce, která bude aplikována nezávisle na každý prvek proudu; je ekvivalentem pixel shaderu. V klasickém programování můžeme uvést analogii smyčky - je aplikována na velký počet Prvky.
  • Chcete-li číst výsledky aplikace jádra na vlákno, musí být vytvořena textura. Na CPU neexistuje žádný ekvivalent, protože má plný přístup k paměti.
  • Řízení umístění v paměti, kde bude proveden zápis (v operacích rozptylu), se provádí pomocí vertex shaderu, protože pixel shader nemůže změnit souřadnice zpracovávaného pixelu.

Jak vidíte, i když vezmeme v úvahu výše uvedené analogie, úkol nevypadá jednoduše. A Brook přišel na pomoc. Tento název odkazuje na rozšíření jazyka C („C with streams“, „C with streams“), jak je nazvali vývojáři ze Stanfordu. V jádru bylo úkolem Brooka skrýt před programátorem všechny komponenty 3D API, což umožnilo prezentovat GPU jako koprocesor pro paralelní výpočty. Za tímto účelem kompilátor Brook zpracoval soubor .br s kódem C++ a příponami a poté vygeneroval kód C++, který byl propojen s knihovnou s podporou různých výstupů (DirectX, OpenGL ARB, OpenGL NV3x, x86).



Pro zvětšení klikněte na obrázek.

Brook má na svém kontě několik zásluh, z nichž první je vyvedení GPGPU ze stínu, aby tuto technologii mohly přijmout masy. I když po oznámení projektu řada IT webů příliš optimisticky hlásila, že vydání Brooka zpochybňuje existenci CPU, která budou brzy nahrazena výkonnějšími GPU. Ale jak vidíme, ani po pěti letech se tak nestalo. Upřímně, nemyslíme si, že se to někdy stane. Na druhou stranu při pohledu na úspěšnou evoluci CPU, která se stále více orientují na paralelismus (více jader, technologie SMT multithreading, rozšiřování SIMD bloků), stejně jako GPU, která jsou naopak stále univerzálnější (podpora výpočtů s plovoucí desetinnou čárkou) jednoduchá přesnost, celočíselné výpočty, podpora výpočtů s dvojnásobnou přesností), zdá se, že GPU a CPU se brzy jednoduše spojí. co se stane potom? Budou GPU absorbovány CPU, jako se to stalo u matematických koprocesorů? Je to docela možné. Intel a AMD dnes pracují na podobných projektech. Ale ještě se může hodně změnit.

Ale vraťme se k našemu tématu. Brookovou výhodou bylo popularizovat koncept GPGPU, což výrazně zjednodušilo přístup ke zdrojům GPU, což umožnilo stále více uživatelům ovládat nový model programování. Na druhou stranu, navzdory všem kvalitám Brooka, zbývala ještě dlouhá cesta k tomu, aby mohly být zdroje GPU použity pro výpočetní techniku.

Jeden z problémů souvisí s na různých úrovních abstrakcí, a také zejména s nadměrným dodatečným zatížením vytvářeným 3D API, které může být docela patrné. Za vážnější problém lze ale považovat problém s kompatibilitou, se kterým vývojáři z Brooku nemohli nic dělat. Mezi výrobci GPU existuje tvrdá konkurence, takže často optimalizují své ovladače. Zatímco takovéto optimalizace jsou většinou dobré pro hráče, mohly by přes noc ukončit Brookovu kompatibilitu. Proto je těžké si představit použití tohoto API v produkčním kódu, který bude někde fungovat. A po dlouhou dobu zůstával Brook výhradou amatérských výzkumníků a programátorů.

Brookův úspěch však stačil k tomu, aby upoutal pozornost ATI a nVidie, které se o takovou iniciativu začaly zajímat, protože by mohla rozšířit trh a otevřít nový důležitý sektor pro společnosti.

Výzkumníci původně zapojení do projektu Brook se rychle spojili s vývojovými týmy v Santa Claře, aby představili globální strategii rozvoje nového trhu. Cílem bylo vytvořit kombinaci hardwaru a software, vhodné pro úlohy GPGPU. Jelikož vývojáři nVidia znají všechna tajemství svých GPU, nemohli se spoléhat na grafické API, ale komunikovat s GPU prostřednictvím ovladače. I když to samozřejmě přichází s vlastními problémy. Vývojový tým CUDA (Compute Unified Device Architecture) tedy vytvořil sadu softwarových vrstev pro práci s GPU.



Pro zvětšení klikněte na obrázek.

Jak můžete vidět na diagramu, CUDA poskytuje dvě API.

  • API na vysoké úrovni: CUDA Runtime API;
  • nízkoúrovňové API: CUDA Driver API.

Protože je vysokoúrovňové rozhraní API implementováno nad nízkoúrovňovým rozhraním API, každé volání funkce na úrovni běhu je rozděleno na více jednoduché instrukce, které jsou zpracovávány rozhraním Driver API. Upozorňujeme, že tato dvě rozhraní API se vzájemně vylučují: programátor může používat jedno nebo druhé rozhraní API, ale nebude možné kombinovat volání funkcí z těchto dvou rozhraní API. Obecně je termín „vysokoúrovňové API“ relativní. Dokonce i Runtime API je takové, že by jej mnozí považovali za nízkoúrovňové; stále však poskytuje funkce, které jsou velmi vhodné pro inicializaci nebo správu kontextu. Ale nečekejte příliš mnoho vysoká úroveň abstrakce - stále musíte mít dobré znalosti o GPU nVidia a o tom, jak fungují.

S rozhraním Driver API je ještě obtížnější pracovat; spuštění zpracování GPU bude vyžadovat více úsilí. Na druhou stranu je nízkoúrovňové API flexibilnější a v případě potřeby poskytuje programátorovi další kontrolu. Dvě rozhraní API jsou schopna pracovat se zdroji OpenGL nebo Direct3D (dnes pouze devátá verze). Přínos této funkce je zřejmý – pomocí CUDA lze vytvářet zdroje (geometrie, procedurální textury atd.), které lze předat grafickému API nebo naopak můžete nechat 3D API posílat výsledky vykreslování do programu CUDA, který zase provede následné zpracování. Existuje mnoho příkladů takových interakcí a výhodou je, že zdroje jsou nadále uloženy v paměti GPU, není nutné je přenášet přes sběrnici PCI Express, která stále zůstává úzkým hrdlem.

Je však třeba poznamenat, že sdílení zdrojů ve videopaměti není vždy ideální a může vést k určitým bolestem hlavy. Například při změně rozlišení nebo barevné hloubky mají přednost grafická data. Pokud tedy potřebujete navýšit prostředky ve vyrovnávací paměti rámců, ovladač to snadno udělá na úkor prostředků CUDA aplikací, které se jednoduše zhroutí s chybou. Samozřejmě ne příliš elegantní, ale taková situace by se neměla stávat příliš často. A protože jsme začali mluvit o nevýhodách: pokud chcete pro aplikace CUDA používat více GPU, musíte nejprve deaktivovat režim SLI, jinak aplikace CUDA budou moci „vidět“ pouze jeden GPU.

Konečně třetí softwarová úroveň je věnována knihovnám – přesněji dvě.

  • CUBLAS, pokud je k dispozici požadované bloky pro výpočty lineární algebry na GPU;
  • CUFFT, který podporuje výpočet Fourierových transformací - algoritmu široce používaného v oblasti zpracování signálů.

Než se ponoříme do CUDA, definujme několik pojmů roztroušených v dokumentaci nVidie. Společnost zvolila velmi specifickou terminologii, na kterou je těžké si zvyknout. Nejprve si toho všimneme vlákno v CUDA nemá stejný význam jako vlákno CPU a také není ekvivalentem vlákna v našich článcích o GPU. GPU závit dovnitř v tomto případě je základní soubor dat, který je třeba zpracovat. Na rozdíl od vláken CPU jsou vlákna CUDA velmi „lehká“, to znamená, že přepínání kontextu mezi dvěma vlákny není operace náročná na zdroje.

Druhý termín, který se často vyskytuje v dokumentaci CUDA, je osnova. Není zde žádný zmatek, protože v ruštině neexistuje žádný analog (pokud nejste fanouškem Start Trek nebo hry Warhammer). Ve skutečnosti tento termín pochází z textilního průmyslu, kde se útková příze protahuje osnovní přízí, která se napíná na tkalcovském stavu. Warp v CUDA je skupina 32 vláken a představuje minimální množství dat zpracovávaných způsobem SIMD v multiprocesorech CUDA.

Ale taková „zrnitost“ není pro programátora vždy vhodná. Proto v CUDA místo přímé práce s osnovami můžete pracovat s bloky, obsahující od 64 do 512 vláken.

Nakonec se tyto bloky spojí mřížka. Výhodou tohoto seskupení je, že počet bloků současně zpracovávaných GPU úzce souvisí s hardwarovými prostředky, jak uvidíme dále. Seskupování bloků do mřížek vám umožňuje zcela abstrahovat toto omezení a aplikovat jádro na více vláken v jediném volání, aniž byste se museli starat o pevné zdroje. To vše mají na svědomí knihovny CUDA. Navíc se takový model dobře měří. Pokud má GPU málo prostředků, bude bloky spouštět postupně. Pokud je počet výpočetních procesorů velký, lze bloky provádět paralelně. To znamená, že na GPU může běžet stejný kód jako vstupní úroveň a na špičkových i budoucích modelech.

V CUDA API je několik dalších termínů, které označují CPU ( hostitel/hostitel) a GPU ( přístroj). Pokud vás tento malý úvod nevyděsil, pak je čas podívat se na CUDA blíže.

Pokud pravidelně čtete Tom's Hardware Guide, pak je vám architektura nejnovějších GPU od nVidie známá. Pokud ne, doporučujeme vám přečíst si článek " nVidia GeForce GTX 260 a 280: nová generace grafických karet„Pokud jde o CUDA, Nvidia prezentuje architekturu trochu jinak a odhaluje některé detaily, které byly dříve skryté.

Jak můžete vidět na obrázku výše, jádro shaderu nVidia se skládá z několika clusterů texturových procesorů (Texture Processor Cluster, TPC). Grafická karta 8800 GTX například používala osm clusterů, 8800 GTS - šest atd. Každý cluster se v podstatě skládá z bloku textury a dvou streamovací multiprocesor. Ty zahrnují začátek pipeline (front end), který čte a dekóduje instrukce a také je posílá k provedení, a konec potrubí (back end), který se skládá z osmi výpočetních zařízení a dvou superfunkčních zařízení. SFU (Super Function Unit), kde jsou instrukce prováděny na principu SIMD, to znamená, že jedna instrukce je aplikována na všechna vlákna ve warpu. nVidia nazývá tento způsob provedení SIMT(jedna instrukce více vláken, jedna instrukce, mnoho vláken). Je důležité si uvědomit, že konec dopravníku pracuje s dvojnásobnou frekvencí než jeho začátek. V praxi to znamená tato část vypadá dvakrát tak „široko“, než ve skutečnosti je (tj. jako 16kanálový blok SIMD namísto osmikanálového). Streamovací multiprocesory fungují následovně: každý cyklus hodin, začátek zřetězení vybere warp připravený k provedení a začne provádět instrukci. Aby se instrukce vztahovala na všech 32 vláken ve warpu, bude konec pipeline vyžadovat čtyři hodinové cykly, ale protože běží s dvojnásobnou frekvencí než začátek, bude vyžadovat pouze dva hodinové cykly (z hlediska začátku potrubí). Aby tedy začátek pipeline nestál nějaký hodinový cyklus naprázdno a hardware byl maximálně zatížen, v ideálním případě můžete každý takt střídat instrukce - klasická instrukce v jednom taktu a instrukce pro SFU v druhém .

Každý multiprocesor má specifickou sadu prostředků, které stojí za pochopení. Existuje malá paměťová oblast tzv "Sdílená paměť", 16 kbajtů na víceprocesor. Nejedná se vůbec o vyrovnávací paměť: programátor ji může použít podle vlastního uvážení. To znamená, že máme něco blízko místního úložiště SPU na procesorech Cell. Tento detail je velmi zajímavý, protože zdůrazňuje, že CUDA je kombinací softwarových a hardwarových technologií. Tato paměťová oblast se nepoužívá pro pixel shadery, na což Nvidia chytře upozorňuje „nelíbí se nám, že spolu pixely mluví“.

Tato oblast paměti otevírá možnost výměny informací mezi vlákny v jednom bloku. Je důležité zdůraznit toto omezení: všechna vlákna v bloku jsou zaručena, že budou spouštěna jedním multiprocesorem. Naproti tomu přiřazení bloků různým multiprocesorům není vůbec specifikováno a dvě vlákna z různých bloků si během provádění nemohou vzájemně vyměňovat informace. To znamená, že použití sdílené paměti není tak snadné. Sdílená paměť je však stále oprávněná s výjimkou případů, kdy se několik vláken pokouší o přístup ke stejné paměťové bance, což způsobuje konflikt. V jiných situacích je přístup ke sdílené paměti stejně rychlý jako přístup k registrům.

Sdílená paměť není jedinou pamětí, ke které mají multiprocesory přístup. Mohou používat videopaměť, ale s menší propustnost a dlouhá zpoždění. Aby se snížila frekvence přístupu k této paměti, nVidia vybavila multiprocesory cache (cca 8 KB na multiprocesor), která ukládá konstanty a textury.

Multiprocesor má 8 192 registrů, které jsou společné pro všechna vlákna všech bloků aktivních na multiprocesoru. Počet aktivních bloků na multiprocesor nesmí překročit osm a počet aktivních warpů je omezen na 24 (768 vláken). Proto 8800 GTX zvládne najednou až 12 288 vláken. Všechna tato omezení stojí za zmínku, protože umožňují optimalizaci algoritmu na základě dostupných zdrojů.

Optimalizace programu CUDA tak spočívá v dosažení optimální rovnováhy mezi počtem bloků a jejich velikostí. Více vláken na blok bude užitečné pro snížení latence paměti, ale také se sníží počet dostupných registrů na vlákno. Navíc blok 512 vláken bude neefektivní, protože na multiprocesoru může být aktivní pouze jeden blok, což vede ke ztrátě 256 vláken. Proto nVidia doporučuje používat bloky o 128 nebo 256 vláknech, což u většiny jader/kernelů poskytuje optimální kompromis mezi sníženou latencí a počtem registrů.

Z hlediska softwaru se CUDA skládá ze sady rozšíření jazyka C, připomínající BrookGPU, a také z několika specifických volání API. Mezi rozšířeními jsou specifikátory typu související s funkcemi a proměnnými. Důležité je zapamatovat si klíčové slovo __globální__, což, když je uvedeno před funkcí, ukazuje, že tato funkce patří k jádru - tato funkce bude volána CPU a bude spuštěna na GPU. Předpona __přístroj__ označuje, že funkce bude spuštěna na GPU (což je mimochodem to, co CUDA nazývá „zařízení“), ale lze ji volat pouze z GPU (jinými slovy z jiné funkce __device__ nebo z funkce __global__) . Nakonec předpona __hostitel__ volitelné, označuje funkci, která je volána CPU a vykonávána CPU - jinými slovy běžná funkce.

S funkcemi __device__ a __global__ souvisí řada omezení: nemohou být rekurzivní (tj. volat samy sebe) a nemohou mít proměnný počet argumentů. A konečně, protože funkce __device__ jsou umístěny v paměťovém prostoru GPU, je logické, že jejich adresu nelze získat. Proměnné mají také řadu kvalifikátorů, které označují oblast paměti, kde budou uloženy. Proměnná s předponou __shared__ znamená, že bude uložen ve sdílené paměti streamovacího multiprocesoru. Volání funkce __global__ je mírně odlišné. Jde o to, že při volání musíte nastavit konfiguraci provádění - konkrétněji velikost mřížky, na kterou bude jádro aplikováno, a také velikost každého bloku. Vezměme si například jádro s následující signaturou.

__global__ void Func(float* parametr);

Bude se jmenovat jako

Func<<< Dg, Db >>>(parametr);

kde Dg je velikost mřížky a Db je velikost bloku. Tyto dvě proměnné odkazují na nový vektorový typ zavedený s CUDA.

CUDA API obsahuje funkce pro práci s pamětí ve VRAM: cudaMalloc pro alokaci paměti, cudaFree pro její uvolnění a cudaMemcpy pro kopírování paměti mezi RAM a VRAM a naopak.

Dokončíme tato recenze Velmi zajímavý způsob kompilace programu CUDA: kompilace se provádí v několika fázích. Nejprve je extrahován kód specifický pro CPU a předán standardnímu kompilátoru. Kód určený pro GPU je nejprve převeden do mezijazyka PTX. Je podobný assembleru a umožňuje vám zkoumat kód a hledat potenciální neefektivitu. A konečně poslední fází je překlad mezijazyka do instrukcí specifických pro GPU a vytvoření binárního souboru.

Po nahlédnutí do dokumentace nVidie jsem v pokušení tento týden vyzkoušet CUDA. Opravdu, co by to mohlo být lepší známky API vytvořením vlastním programem? Právě tehdy by měla vyplout na povrch většina problémů, i když na papíře vše vypadá perfektně. Praxe také nejlépe ukáže, jak dobře jste pochopili všechny principy nastíněné v dokumentaci CUDA.

Zapojit se do takového projektu je docela snadné. Dnes je ke stažení velké množství bezplatných, ale kvalitních nástrojů. Pro náš test jsme použili Visual C++ Express 2005, který má vše, co potřebujeme. Nejtěžší bylo najít program, jehož portování na GPU by netrvalo týdny, ale zároveň by byl dostatečně zajímavý, aby naše úsilí nepřišlo nazmar. Nakonec jsme zvolili kus kódu, který vezme výškovou mapu a vypočítá odpovídající normální mapu. O této funkci se nebudeme podrobně věnovat, protože v tomto článku je sotva zajímavá. Abychom byli struční, program se zabývá zakřivením oblastí: pro každý pixel výchozího obrázku vložíme matici, která pomocí více či méně složitého vzorce určuje barvu výsledného pixelu ve vygenerovaném obrázku ze sousedních pixelů. Výhodou této funkce je velmi snadná paralelizace, takže tento test dokonale ukazuje možnosti CUDA.


Další výhodou je, že již máme implementaci CPU, takže výsledek můžeme porovnat s verzí CUDA – aniž bychom znovu vynalézali kolo.

Ještě jednou zopakujme, že účelem testu bylo seznámit se s utilitami CUDA SDK, nikoli porovnávat verze pro CPU a GPU. Protože to byl náš první pokus o vytvoření programu CUDA, neočekávali jsme vysoký výkon. Protože tato část kódu není kritická, verze pro CPU nebyla optimalizována, takže přímé srovnání výsledků je sotva zajímavé.

Výkon

Měřili jsme však dobu provádění, abychom zjistili, zda existuje výhoda použití CUDA i při nejhrubší implementaci, nebo zda bychom potřebovali dlouhé a únavné cvičení, abychom získali nějaké výhody při používání GPU. Testovací stroj byl převzat z naší vývojové laboratoře - notebook s Jádrový procesor 2 Duo T5450 a grafická karta GeForce 8600M GT se systémem Vista. Nejde o superpočítač, ale výsledky jsou velmi zajímavé, protože test není „šitý na míru“ GPU. Vždy je hezké vidět, že Nvidia ukazuje obrovské zisky na systémech s monstrózními GPU a velkou šířkou pásma, ale v praxi mnoho ze 70 milionů GPU s podporou CUDA na dnešním trhu PC není zdaleka tak výkonných, a proto náš test stále platí. .

Pro obrázek o velikosti 2048 x 2048 pixelů jsme získali následující výsledky.

  • CPU 1 vlákno: 1 419 ms;
  • CPU 2 vlákna: 749 ms;
  • CPU 4 vlákna: 593 ms
  • GPU (8600M GT) bloky 256 vláken: 109 ms;
  • GPU (8600M GT) bloky 128 vláken: 94 ms;
  • GPU (8800 GTX) bloky 128 vláken/256 vláken: 31 ms.

Z výsledků lze vyvodit několik závěrů. Začněme tím, že i přes řeči o zjevné lenosti programátorů jsme upravili výchozí verzi CPU pro více vláken. Jak jsme již zmínili, kód je pro tuto situaci ideální – vše, co je potřeba, je rozdělit počáteční obrázek do tolika zón, kolik je vláken. Vezměte prosím na vědomí, že přechod z jednoho vlákna na dva na našem dvoujádrovém CPU vedl k téměř lineárnímu zrychlení, což také naznačuje paralelní povahu testovacího programu. Zcela nečekaně se rychlejší ukázala i verze se čtyřmi vlákny, i když na našem procesoru je to velmi zvláštní – naopak se dal očekávat pokles efektivity kvůli režii správy dalších vláken. Jak lze tento výsledek vysvětlit? Těžko říct, ale může za to plánovač vláken Windows; v každém případě výsledek opakujeme. S menšími texturami (512x512) nebyl zisk z vláken tak výrazný (asi 35 % vs. 100 %) a chování čtyřvláknové verze bylo logičtější, bez nárůstu oproti dvouvláknové verzi. GPU bylo stále rychlejší, ale ne tak znatelně (8600M GT bylo třikrát rychlejší než dvouvláknová verze).



Pro zvětšení klikněte na obrázek.

Druhým významným postřehem je, že i nejpomalejší implementace GPU byla téměř šestkrát rychlejší než nejvýkonnější verze CPU. U prvního programu a neoptimalizované verze algoritmu je výsledek velmi povzbudivý. Vezměte prosím na vědomí, že jsme obdrželi znatelně nejlepší výsledek na malých blocích, i když intuice může naznačovat opak. Vysvětlení je jednoduché – náš program používá 14 registrů na vlákno a při 256vláknových blocích je potřeba 3 584 registrů na blok a k plnému zatížení procesoru je potřeba 768 vláken, jak jsme si ukázali. V našem případě to činí tři bloky nebo 10 572 registrů. Ale multiprocesor má pouze 8 192 registrů, takže dokáže udržet aktivní pouze dva bloky. Naproti tomu u bloků o 128 vláknech potřebujeme 1 792 registrů na blok; Pokud je 8 192 děleno 1 792 a zaokrouhleno na nejbližší celé číslo, dostaneme čtyři bloky. V praxi bude počet vláken stejný (512 na multiprocesor, i když teoreticky jich je pro plné zatížení potřeba 768), ale zvýšení počtu bloků dává GPU výhodu flexibility v přístupu k paměti – když dojde k operaci s velké zpoždění, můžete začít provádět instrukce z jiného bloku, zatímco čekáte na obdržení výsledků. Čtyři bloky jasně snižují latenci, zejména proto, že náš program používá více přístupů do paměti.

Analýza

Nakonec, navzdory tomu, co jsme řekli výše, jsme neodolali pokušení a spustili jsme program na 8800 GTX, která byla třikrát rychlejší než 8600, bez ohledu na velikost bloku. Možná si myslíte, že v praxi na odpovídajících architekturách bude výsledek čtyři a vícekrát vyšší: 128 ALU/shader procesorů oproti 32 a vyšší takt (1,35 GHz versus 950 MHz), ale tak to nefungovalo. S největší pravděpodobností byl limitujícím faktorem přístup do paměti. Přesněji řečeno, k počátečnímu obrázku se přistupuje jako k multidimenzionálnímu poli CUDA – poměrně složitý termín pro to, co není nic jiného než textura. Ale existuje několik výhod.

  • přístupy těží z mezipaměti textur;
  • používáme režim zalamování, ve kterém není potřeba zpracovávat hranice obrazu, na rozdíl od verze CPU.

Kromě toho můžeme těžit z „bezplatného“ filtrování s normalizovaným adresováním mezi místo a , ale v našem případě to pravděpodobně nebude užitečné. Jak víte, 8600 má 16 texturových jednotek ve srovnání s 32 u 8800 GTX. Proto je poměr mezi těmito dvěma architekturami pouze dva ku jedné. Přidejte k tomu rozdíl ve frekvencích a dostaneme poměr (32 x 0,575) / (16 x 0,475) = 2,4 - blízko poměru "tři ku jedné", který jsme skutečně dostali. Tato teorie také vysvětluje, proč se velikost bloku na G80 příliš nemění, protože ALU stále spočívá na blocích textury.



Pro zvětšení klikněte na obrázek.

Kromě slibných výsledků dopadlo naše první seznámení s CUDA velmi dobře, vzhledem k zvoleným ne zrovna nejpříznivějším podmínkám. Vývoj na notebooku Vista znamená, že budete muset používat CUDA SDK 2.0, stále v beta verzi, s ovladačem 174.55, který je také beta. Přesto nemůžeme hlásit žádná nepříjemná překvapení – pouze prvotní chyby při prvním ladění, kdy se náš stále velmi zabugovaný program snažil adresovat paměť mimo přidělený prostor.

Monitor začal divoce blikat, pak obrazovka zčernala... dokud Vista nespustila službu opravy ovladačů a vše bylo v pořádku. Ale stále je to poněkud překvapivé, pokud jste zvyklí vidět typickou chybu Segmentation Fault ve standardních programech, jako je ten náš. Na závěr malá výtka směrem k nVidii: v celé dokumentaci dostupné pro CUDA není žádný malý průvodce, který by vám krok za krokem řekl, jak nastavit vývojové prostředí pro Visual Studio. Ve skutečnosti to není velký problém, protože SDK má úplnou sadu příkladů, které si můžete prostudovat, abyste pochopili rámec pro aplikace CUDA, ale průvodce pro začátečníky by byl užitečný.



Pro zvětšení klikněte na obrázek.

Nvidia představila CUDA s vydáním GeForce 8800. A v té době se zdály sliby velmi lákavé, ale naše nadšení jsme zadrželi, dokud jsme to skutečně neotestovali. Ve skutečnosti to v té době vypadalo spíše jako označení území, aby zůstalo na vlně GPGPU. Bez dostupné sady SDK je těžké říci, že to není jen další marketingový blábol, který nebude fungovat. Není to poprvé, co byla dobrá iniciativa ohlášena příliš brzy a kvůli nedostatečné podpoře tehdy nespatřila světlo světa – zvláště v tak konkurenčním odvětví. Nyní, rok a půl po oznámení, můžeme s jistotou říci, že nVidia dodržela slovo.

SDK se rychle objevilo v beta verzi na začátku roku 2007 a od té doby bylo rychle aktualizováno, což dokazuje důležitost tohoto projektu pro nVidia. Dnes se CUDA vyvíjí velmi pěkně: SDK je již k dispozici v beta verzi 2.0 pro hlavní operační systémy (Windows XP a Vista, Linux a také 1.1 pro Mac OS X) a nVidia věnovala celou jednu sekci webu pro vývojáři.

Na profesionálnější úrovni dopadl dojem z prvních kroků s CUDA velmi pozitivní. I když jste obeznámeni s architekturou GPU, můžete na to snadno přijít. Když API vypadá na první pohled jasně, okamžitě začnete věřit, že získáte přesvědčivé výsledky. Neztratí se ale výpočetní čas četnými přenosy z CPU do GPU? A jak používat tyto tisíce vláken prakticky bez synchronizačního primitiva? Začali jsme experimentovat se všemi těmito obavami. Ale rychle se rozplynuly, když se první verze našeho algoritmu, i když velmi triviální, ukázala být výrazně rychlejší než na CPU.

CUDA tedy není kouzelná hůlka pro výzkumníky, kteří chtějí přesvědčit vedení univerzity, aby jim koupilo GeForce. CUDA je již zcela dostupná technologie, kterou může používat každý programátor se znalostí C, pokud je ochoten věnovat čas a úsilí, aby si zvykl na nové programovací paradigma. Toto úsilí nebude zbytečné, pokud jsou vaše algoritmy dobře paralelizované. Rádi bychom také poděkovali společnosti nVidia za poskytnutí kompletní a kvalitní dokumentace, kde najdou odpovědi začínající programátoři CUDA.

Co potřebuje CUDA, aby se stalo rozpoznatelným API? Jedním slovem: přenositelnost. Víme, že budoucnost IT spočívá v paralelním počítání – dnes se již na takové změny všichni připravují a všechny iniciativy, softwarové i hardwarové, míří tímto směrem. Nicméně v tuto chvíli, když se podíváme na vývoj paradigmat, jsme stále v počáteční fázi: vytváříme vlákna ručně a snažíme se naplánovat přístup k sdílené zdroje; S tím vším se dá nějak vypořádat, pokud se počet jader dá spočítat na prstech jedné ruky. Jenže za pár let, kdy se počet procesorů bude počítat ve stovkách, tato možnost již nebude. S vydáním CUDA udělala nVidia první krok k vyřešení tohoto problému – ale samozřejmě toto rozhodnutí Vhodné pouze pro GPU od této společnosti a i tak ne pro každého. Programy CUDA dnes umí spustit pouze GF8 a 9 (a jejich deriváty Quadro/Tesla). A samozřejmě nová řada 260/280.



Pro zvětšení klikněte na obrázek.

Nvidia se může chlubit, že po celém světě prodala 70 milionů GPU kompatibilních s CUDA, ale stále to nestačí na to, aby se stala de facto standardem. S přihlédnutím k tomu, že konkurenti nesedí nečinně. AMD nabízí vlastní SDK (Stream Computing) a Intel oznámil řešení (Ct), i když zatím není k dispozici. Blíží se válka standardů a na trhu zjevně nezbude místo pro tři konkurenty, dokud další hráč jako Microsoft nepřijde se společným API, které vývojářům jistě usnadní život.

Proto má nVidia mnoho problémů se schválením CUDA. Přestože technologicky stojíme nepochybně před úspěšným řešením, zbývá ještě přesvědčit vývojáře o jeho perspektivách – a to nebude jednoduché. Nicméně, soudě podle mnoha nedávných oznámení a novinek týkajících se API, budoucnost nevypadá chmurně.

– sada nízkoúrovňových softwarových rozhraní ( API) pro vytváření her a dalších vysoce výkonných multimediálních aplikací. Zahrnuje podporu vysokého výkonu 2D- A 3D-grafika, zvuk a vstupní zařízení.

Direct3D (D3D) – rozhraní pro zobrazování trojrozměrných primitivů(geometrická tělesa). Obsažen v .

OpenGL(z angličtiny Otevřete grafickou knihovnu, doslova - otevřená grafická knihovna) je specifikace, která definuje meziplatformní programovací rozhraní nezávislé na programovacím jazyce pro psaní aplikací využívajících dvourozměrnou a trojrozměrnou počítačovou grafiku. Obsahuje více než 250 funkcí pro kreslení složitých 3D scén z jednoduchých primitiv. Používá se k vytváření videoher, virtuální reality a vizualizací ve vědeckém výzkumu. Na platformě Okna soutěží s .

OpenCL(z angličtiny Otevřete Computing Language, doslova – otevřený jazyk výpočtů) – rámec(softwarový systémový framework) pro psaní počítačových programů souvisejících s paralelními výpočty na různých grafikách ( GPU) A ( ). Do rámce OpenCL obsahuje programovací jazyk a rozhraní pro programování aplikací ( API). OpenCL poskytuje paralelismus na úrovni instrukcí a na úrovni dat a je implementací této techniky GPGPU.

GPGPU(zkráceně z angličtiny) G rafické zpracovatelské jednotky pro všeobecné použití, doslova - GPU obecný účel) je technika pro použití grafického procesoru (GPU) nebo grafické karty pro obecné výpočty, které obvykle provádí počítač.

Shader(Angličtina) shader) – program pro konstrukci stínů na syntetizovaných obrázcích, používaný v trojrozměrné grafice k určení konečných parametrů objektu nebo obrázku. Typicky zahrnuje libovolně složité popisy absorpce a rozptylu světla, mapování textur, odraz a lom, stínování, posun povrchu a efekty následného zpracování. Složité povrchy lze vizualizovat pomocí jednoduchých geometrických tvarů.

Vykreslování(Angličtina) vykreslování) – vizualizace, in počítačová grafika proces získání obrázku z modelu pomocí softwaru.

SDK(zkráceně z angličtiny) Software Development Kit) – sada nástrojů pro vývoj softwaru.

procesor(zkráceně z angličtiny) Centrální procesorová jednotka, doslova – centrální/hlavní/hlavní výpočetní zařízení) – centrální (mikro);zařízení, které provádí strojové instrukce; Část Hardware, zodpovědný za provádění výpočetních operací (specifikováno operační systém a aplikační software) a koordinace provozu všech zařízení.

GPU(zkráceně z angličtiny) Grafická procesorová jednotka, doslova – grafické výpočetní zařízení) – grafický procesor; samostatné zařízení nebo herní konzole, která provádí grafické vykreslování (vizualizace). Moderní GPU jsou velmi efektivní při zpracování a zobrazení počítačové grafiky realistickým způsobem. Grafický procesor v moderních grafických adaptérech se používá jako 3D grafický akcelerátor, ale v některých případech jej lze použít i pro výpočty ( GPGPU).

Problémy procesor

Ke zvýšení výkonu tradičních docházelo dlouhou dobu především důsledným zvyšováním taktovací frekvence (cca 80 % výkonu bylo určeno taktovací frekvencí) se současným nárůstem počtu tranzistorů na jednom čipu. . Další zvýšení taktovací frekvence (při taktovací frekvenci více než 3,8 GHz se čipy jednoduše přehřívají!) však naráží na řadu zásadních fyzických překážek (protože technologický proces se téměř přiblížil velikosti atomu: , a velikost atomu křemíku je přibližně 0,543 nm):

Za prvé, jak se zmenšuje velikost krystalu a zvyšuje se frekvence hodin, zvyšuje se svodový proud tranzistorů. To vede ke zvýšené spotřebě energie a zvýšeným emisím tepla;

Za druhé, výhody vyšší rychlosti hodin jsou částečně negovány latencí přístupu do paměti, protože doba přístupu do paměti nedrží krok s rostoucí rychlostí hodin;

Zatřetí, pro některé aplikace se tradiční sériové architektury stávají neefektivními, protože se takty zvyšují kvůli takzvanému „von Neumannovu úzkému hrdlu“, což je omezení výkonu vyplývající ze sekvenčního výpočetního toku. Současně se zvětšují zpoždění při přenosu odporově-kapacitního signálu, což je další úzké hrdlo spojené se zvýšením taktovací frekvence.

Rozvoj GPU

Paralelně s tím probíhal (a probíhá!) vývoj GPU:

listopad 2008 – Intel představila řadu 4jádrových Intel Core i7, které jsou založeny na nové generaci mikroarchitektury Nehalem. Procesory pracují na taktovací frekvenci 2,6-3,2 GHz. Vyrobeno pomocí 45nm procesní technologie.

Prosinec 2008 – zahájeny dodávky 4jádrových AMD Phenom II 940(krycí jméno - Deneb). Pracuje na frekvenci 3 GHz, vyrobeno pomocí 45nm procesní technologie.

květen 2009 – spol AMD představila verzi GPU Grafická karta ATI Radeon HD 4890 S hodinová frekvence jádro zvýšeno z 850 MHz na 1 GHz. Toto je první grafický procesor běží na 1 GHz. Výpočetní výkon čipu se díky zvýšení frekvence zvýšil z 1,36 na 1,6 teraflopů. Procesor obsahuje 800 (!) výpočetních jader a podporuje videopaměť GDDR5, DirectX 10.1, ATI CrossFireX a všechny další související technologie moderní modely grafické karty Čip je vyroben na bázi 55nm technologie.

Hlavní rozdíly GPU

Charakteristické rysy GPU(ve srovnání s ) jsou:

– architektura maximálně zaměřená na zvýšení rychlosti výpočtu textur a složitých grafických objektů;

– typický špičkový výkon GPU mnohem vyšší než to ;

– díky specializované dopravníkové architektuře, GPU mnohem efektivnější při zpracování grafické informace, jak .

"Krize žánru"

"Žánrová krize" pro dozrály v roce 2005 – tehdy se objevily. Ale i přes vývoj technologií, zvýšení produktivity konvenčních znatelně klesla. Zároveň výkon GPU stále roste. Takže v roce 2003 tato revoluční myšlenka vykrystalizovala - využijte výpočetní výkon grafiky pro své potřeby. GPU se stále více používají pro „negrafické“ výpočty (fyzikální simulace, zpracování signálů, výpočetní matematika/geometrie, operace s databázemi, výpočetní biologie, výpočetní ekonomika, počítačové vidění atd.).

Hlavním problémem bylo, že neexistovalo standardní programovací rozhraní GPU. Vývojáři použili OpenGL nebo Direct3D, ale bylo to velmi pohodlné. Korporace NVIDIA(jeden z největších výrobců grafických, mediálních a komunikačních procesorů a také bezdrátových mediálních procesorů; založena v roce 1993) začala vyvíjet jednotný a pohodlný standard - a představila technologii CUDA.

Jak to začalo

2006 – NVIDIA demonstruje CUDA™; začátek revoluce ve výpočetní technice GPU.

2007 – NVIDIA uvolňuje architekturu CUDA(originální verze CUDA SDK byla podána 15. února 2007); nominace" Nejlepší nový produkt» z časopisu Populární věda a "Volba čtenářů" z publikace HPCWire.

2008 – technologie NVIDIA CUDA zvítězil v kategorii „Technická dokonalost“. PC magazín.

Co se stalo CUDA

CUDA(zkráceně z angličtiny) Compute Unified Device Architecture, doslova - jednotná výpočetní architektura zařízení) - architektura (soubor softwaru a hardwaru), která umožňuje vyrábět na GPU obecné výpočty, zatímco GPU ve skutečnosti funguje jako výkonný koprocesor.

Technika NVIDIA CUDA™ je jediné vývojové prostředí v programovacím jazyce C, která umožňuje vývojářům vytvářet software, který řeší složité výpočetní problémy v kratším čase, a to díky výpočetnímu výkonu GPU. Ve světě už pracují miliony lidí GPU s podporou CUDA a tisíce programátorů již používají (bezplatné!) nástroje CUDA k urychlení aplikací a řešení nejsložitějších úkolů náročných na zdroje – od kódování videa a zvuku po průzkum ropy a zemního plynu, modelování produktů, lékařské zobrazování a vědecký výzkum.

CUDA dává vývojáři možnost podle vlastního uvážení organizovat přístup k sadě instrukcí grafického akcelerátoru a spravovat jeho paměť a organizovat na něm složité paralelní výpočty. Podpora grafického akcelerátoru CUDA se stává výkonnou programovatelnou otevřenou architekturou, podobnou dnešní . To vše poskytuje vývojářům nízkoúrovňový, distribuovaný a vysokorychlostní přístup k hardwaru CUDA nezbytný základ pro budování seriózních nástrojů na vysoké úrovni, jako jsou kompilátory, debuggery, matematické knihovny a softwarové platformy.

Uralsky, přední technologický specialista NVIDIA, srovnávání GPU A , říká toto: „ - Tohle je SUV. Jezdí vždy a všude, ale ne moc rychle. A GPU- Tohle je sportovní auto. Na špatné silnici prostě nikam nedojede, ale dá mu pořádný povrch a ukáže svou rychlost, o které se SUV ani nesnilo!…

Technologické schopnosti CUDA

Jádra CUDAsymbol skalární výpočetní jednotky ve video čipech NVidia, počínaje G 80 (GeForce 8 xxx, Tesla C-D-S870, FX4/5600 , 360 mil). Samotné čipy jsou deriváty architektury. Mimochodem, protože spol NVidia se tak ochotně ujala vývoje vlastních procesorů Série Tegra, také na základě RISC architektura. S prací s těmito architekturami mám bohaté zkušenosti.

CUDA jádro obsahuje jeden jeden vektor A jeden skalár jednotky, které provádějí jednu vektorovou a jednu skalární operaci za cyklus hodin, přenášejí výpočty do jiného multiprocesoru nebo do jiného pro další zpracování. Pole stovek a tisíců takových jader představuje významný výpočetní výkon a může provádět různé úkoly v závislosti na požadavcích za předpokladu, že existuje určitý podpůrný software. aplikace lze měnit: dekódování video streamu, 2D/3D grafická akcelerace, cloud computing, specializované matematické analýzy atd.

Dost často kombinované profesionální karty NVidia Tesla A NVidia Quadro, jsou páteří moderních superpočítačů.

CUDA— jádra od té doby neprošla žádnými významnými změnami G 80, ale jejich počet se zvyšuje (spolu s dalšími bloky - ROP, Texturní jednotky atd.) a účinnost paralelních interakcí mezi sebou (moduly jsou vylepšeny Giga vlákno).

Např:

GeForce

Jádra GTX 460 – 336 CUDA
Jádra GTX 580 – 512 CUDA
8800GTX – 128 CUDA jader

Z počtu stream procesorů ( CUDA), výkon ve výpočtech shaderu se zvyšuje téměř úměrně (s rovnoměrným nárůstem počtu dalších prvků).

Počínaje čipem GK110 (NVidia GeForce GTX 680) — CUDA jádra již nemají dvojnásobnou frekvenci, ale společnou se všemi ostatními čipovými bloky. Místo toho se jejich počet přibližně zvýšil třikrát oproti předchozí generaci G110.

Podle Darwinovy ​​evoluční teorie první opice (pokud
abych byl přesný - homo antecessor, lidský předchůdce) se později změnil v
v nás. Vícetunová počítačová centra s tisíci nebo více rádiovými trubicemi,
zabírající celé místnosti byly nahrazeny půlkilogramovými notebooky, které mimochodem
nebude výkonově horší než první. Předpotopní psací stroje se staly
při tisku čehokoli a na cokoli (i na lidské tělo)
multifunkční zařízení. Obři procesorů se najednou rozhodli zazdít
grafické jádro v "kámeni". A grafické karty začaly nejen zobrazovat obrázek
přijatelnou FPS a kvalitu grafiky, ale také provádět všechny druhy výpočtů. Ano
stále jak vyrábět! Bude diskutována technologie vícevláknových výpočtů pomocí GPU.

Proč GPU?

Zajímalo by mě, proč se rozhodli přenést veškerý výpočetní výkon na grafiku
adaptér? Jak vidíte, procesory jsou stále v módě a je nepravděpodobné, že se vzdají svého tepla
místo. Ale GPU má v rukávu pár es, spolu s žolíkem a pár rukávů
dost. Moderní centrální procesor je navržen tak, aby dosahoval maxima
výkon při zpracování celých čísel a dat s pohyblivou řádovou čárkou
čárkou, aniž by se zvlášť staral o paralelní zpracování informací. Zároveň
architektura grafické karty vám umožní rychle a bez problémů „paralelizovat“
zpracování dat. Na jedné straně se počítají polygony (kvůli 3D dopravníku),
na druhou stranu zpracování pixelové textury. Je jasné, že existuje „harmonický
rozpadu zátěže v jádru karty. Navíc výkon paměti a video procesoru
optimálnější než kombinace „RAM-cache-processor“. V okamžiku, kdy datová jednotka
na grafické kartě začne být zpracováván jedním GPU stream procesorem, druhým
Jednotka je načtena paralelně do jiné a v zásadě je snadné toho dosáhnout
zatížení GPU srovnatelné s šířkou pásma sběrnice,
k tomu však musí být dopravníky zatíženy rovnoměrně, bez
jakékoli podmíněné přechody a větvení. Centrální procesor, na základě jeho
všestrannost vyžaduje plnou mezipaměť pro potřeby zpracování
informace.

Vědci přemýšleli o práci GPU v paralelním počítání a
matematiky a přišel s teorií, které jsou mnohé vědecké výpočty v mnohém podobné
Zpracování 3D grafiky. Mnoho odborníků se domnívá, že základním faktorem v
rozvoj GPGPU (General Purpose computing na GPU – univerzální
výpočty pomocí grafické karty
) byl vznik projektu Brook GPU v roce 2003.

Tvůrci projektu ze Stanfordské univerzity museli vyřešit obtížný úkol
problém: hardware a software nutit grafický adaptér vyrábět
různé výpočty. A uspěli. Pomocí obecného jazyka C,
Američtí vědci přiměli GPU pracovat jako procesor, upravený pro
paralelní zpracování. Po Brookovi se objevila řada projektů na výpočty VGA,
jako je knihovna Accelerator, knihovna Brahma, systém
GPU++ metaprogramování a další.

CUDA!

Předtucha vyhlídek rozvoje vynucená AMD A NVIDIA
držet se Brook GPU jako pitbull. Když pomineme marketingovou politiku, tak
Správnou implementací všeho se můžete prosadit nejen v grafickém sektoru
trhu, ale také v oblasti výpočetní techniky (podívejte se na speciální počítačové karty a
servery Tesla se stovkami multiprocesorů), které nahrazují obvyklé CPU.

Přirozeně, že „páni FPS“ se rozešli u kamene úrazu, každý po svém.
cesta, ale základní princip zůstal nezměněn - provádět výpočty
pomocí GPU. A nyní se blíže podíváme na „zelenou“ technologii - CUDA
(Compute Unified Device Architecture).

Úkolem naší „hrdinky“ je poskytnout API, dvě najednou.
První je vysokoúrovňová, CUDA Runtime, která představuje funkce, které
jsou rozděleny do jednodušších úrovní a předány nižšímu API - CUDA Driver. Tak
že fráze „na vysoké úrovni“ je úsek, který lze na proces aplikovat. Veškerá sůl je
přesně v ovladači a laskavě vytvořené knihovny vám jej pomohou získat
vývojáři NVIDIA: CUBLAS (nástroje pro matematické výpočty) a
FFT (výpočet pomocí Fourierova algoritmu). No, přejděme k praktickému
části materiálu.

Terminologie CUDA

NVIDIA pracuje s velmi unikátními definicemi pro CUDA API. Ony
se liší od definic používaných pro práci s centrálním procesorem.

Vlákno– soubor údajů, které je třeba zpracovat (ne
vyžaduje velké zdroje na zpracování).

Warp– skupina 32 vláken. Údaje se pouze zpracovávají
warp, proto je warp minimální množství dat.

Blok– sada toků (od 64 do 512) nebo sada
warps (od 2 do 16).

Mřížka je sbírka bloků. Toto rozdělení dat
slouží výhradně ke zlepšení výkonu. Pokud tedy číslo
multiprocesorů je velký, pak budou bloky vykonávány paralelně. Pokud s
žádné štěstí s kartou (vývojáři doporučují použít
adaptér ne nižší než GeForce 8800 GTS 320 MB), pak budou datové bloky zpracovány
postupně.

NVIDIA také zavádí pojmy jako např jádro, hostitel
A přístroj.

Pracovali!

Pro plnou práci s CUDA potřebujete:

1. Znát strukturu shader jader GPU, od podstaty programování
spočívá v rovnoměrném rozložení zátěže mezi ně.
2. Umět programovat v prostředí C s přihlédnutím k některým aspektům.

Vývojáři NVIDIA několikrát odhalil „vnitřky“ grafické karty
jinak, než jsme zvyklí vidět. Takže si chtě nechtě budete muset všechno nastudovat
jemnosti architektury. Podívejme se na strukturu legendárního „kámenu“ G80 GeForce 8800
GTX
.

Jádro shaderu se skládá z osmi clusterů TPC (Texture Processor Cluster).
texturové procesory (takže GeForce GTX 280- 15 jader, 8800 GTS
je jich šest 8600 – čtyři atd.). Ty se zase skládají ze dvou
streamovací multiprocesory (dále jen SM). SM (všechny
16) sestává z frontendu (řeší problémy čtení a dekódování instrukcí) a
back-end (konečný výstup instrukcí) potrubí, stejně jako osm skalárních SP (shader
procesor) a dvě SFU (super funkční jednotky). Za každý úder (jednotka
time) frontend vybere warp a zpracuje ho. Aby veškerý warp tekl
(připomínám, je jich 32) zpracováno, na konci dopravníku jsou potřeba 32/8 = 4 cykly.

Každý multiprocesor má takzvanou sdílenou paměť.
Jeho velikost je 16 kilobajtů a poskytuje programátorovi naprostou svobodu
akce. Distribuujte dle libosti :). Sdílená paměť zajišťuje komunikaci mezi vlákny
jeden blok a není určen pro práci s pixel shadery.

SM mají také přístup k GDDR. K tomu dostali každý 8 kilobajtů.
mezipaměť, která ukládá všechny nejdůležitější věci pro práci (například výpočetní techniku
konstanty).

Multiprocesor má 8192 registrů. Počet aktivních bloků nemůže být
více než osm a počet warpů není větší než 768/32 = 24. Z toho je zřejmé, že G80
dokáže zpracovat maximálně 32*16*24 = 12288 vláken za jednotku času. Nemůžeš si pomoct
vzít tato čísla v úvahu při optimalizaci programu v budoucnu (na jednom měřítku
– velikost bloku, na druhé straně – počet vláken). Roli může hrát vyváženost parametrů
proto důležitou roli v budoucnu NVIDIA doporučuje používat bloky
se 128 nebo 256 vlákny. Blok 512 vláken je neefektivní, protože má
zvýšené zpoždění. S přihlédnutím ke všem jemnostem struktury grafické karty GPU plus
dobré programovací dovednosti, můžete vytvářet velmi produktivní
nástroj pro paralelní výpočty. Mimochodem, o programování...

Programování

Pro „kreativitu“ s CUDA potřebujete Grafická karta GeForce není nižší
epizoda osmá
. S

z oficiálních stránek si musíte stáhnout tři softwarové balíčky: ovladač
Podpora CUDA (každý OS má svůj), samotný balíček CUDA SDK (druhý
beta verze) a další knihovny (CUDA toolkit). Technologie podporuje
operační systémy Windows (XP a Vista), Linux a Mac OS X. Ke studiu I
zvolil Vista Ultimate Edition x64 (při pohledu dopředu řeknu, že se systém choval
Prostě skvělé). V době psaní těchto řádků to bylo pro práci relevantní
ForceWare ovladač 177.35. Používá se jako sada nástrojů
Softwarový balík Borland C++ 6 Builder (ačkoli jakékoli prostředí, které pracuje s
jazyk C).

K osobě kdo zná jazyk, bude snadné si zvyknout na nové prostředí. Vše, co je požadováno, je
zapamatovat si základní parametry. Klíčové slovo _global_ (umístěné před funkcí)
označuje, že funkce patří k jádru. Zavolá ji centrála
procesor a veškerá práce bude probíhat na GPU. Volání _global_ vyžaduje více
konkrétní detaily, konkrétně velikost sítě, velikost bloku a jaké bude jádro
aplikovaný. Například řádek _global_ void saxpy_parallel<<>>, kde X –
velikost mřížky a Y je velikost bloku, určuje tyto parametry.

Symbol _device_ znamená, že funkci bude volat grafické jádro, známé také jako
bude dodržovat všechny pokyny. Tato funkce je umístěna v paměti multiprocesoru,
proto je nemožné získat její adresu. Předpona _host_ znamená, že volání
a zpracování bude probíhat pouze za účasti CPU. Je třeba vzít v úvahu, že _globální_ a
_zařízení_ si nemohou volat navzájem a nemohou volat sami.

Jazyk pro CUDA má také řadu funkcí pro práci s videopamětí: cudafree
(uvolnění paměti mezi GDDR a RAM), cudamemcpy a cudamemcpy2D (kopírování
paměti mezi GDDR a RAM) a cudamalloc (alokace paměti).

Všechny programové kódy jsou kompilovány pomocí CUDA API. Nejprve se bere
kód určený výhradně pro centrální procesor a je vystaven
standardní kompilace a další kód určený pro grafický adaptér
přepsáno do středního jazyka PTX (podobně jako assembler).
identifikace možné chyby. Po všech těchto „tancích“ finále
překlad (překlad) příkazů do jazyka srozumitelného pro GPU/CPU.

Studijní sada

Téměř všechny aspekty programování jsou popsány v dokumentaci
spolu s ovladačem a dvěma aplikacemi a také na webu vývojářů. Velikost
článek k jejich popisu nestačí (zájemce by měl přiložit
trochu úsilí a prostudujte si materiál sami).

CUDA SDK Browser byl vyvinut speciálně pro začátečníky. Kdokoli může
pocítit sílu paralelního počítání na vlastní kůži (nejlepší test pro
stabilita – příklady fungují bez artefaktů nebo pádů). Aplikace má
velký počet orientačních miniprogramů (61 „testů“). Pro každou zkušenost existuje
podrobnou dokumentaci programový kód plus soubory PDF. Okamžitě je zřejmé, že lidé
přítomní se svými výtvory v prohlížeči odvádějí seriózní práci.
Můžete také porovnat rychlost procesoru a grafické karty během zpracování
data. Například skenování vícerozměrných polí pomocí grafické karty GeForce 8800
GT
Vyrobí 512 MB s blokem s 256 vlákny za 0,17109 milisekund.
Technologie nerozpoznává tandemy SLI, takže pokud máte duo nebo trio,
před zahájením práce deaktivujte funkci „párování“, jinak CUDA uvidí pouze jednu
přístroj Dvoujádrový AMD Athlon 64 X2(frekvence jádra 3000 MHz) stejná zkušenost
projde za 2,761528 milisekund. Ukazuje se, že G92 je více než 16krát
rychlejší než kámen AMD! Jak vidíte, nejde zdaleka o extrémní systém
v tandemu s operačním systémem nemilovaným masami ukazuje dobré
Výsledek.

Kromě prohlížeče existuje řada programů užitečných pro společnost. Adobe
přizpůsobila své produkty nová technologie. Nyní je Photoshop CS4 v plném rozsahu
nejméně využívá zdroje grafické adaptéry(musíte si stáhnout speciální
zapojit). S programy jako Badaboom media converter a RapiHD můžete
dekódovat video do formátu MPEG-2. Dobré pro zpracování zvuku
udělám bezplatný nástroj Accelero. množství softwaru přizpůsobeného pro CUDA API,
nepochybně poroste.

A v této době...

Mezitím čtete tento materiál, dříči z zpracovatelských koncernů
vyvíjejí vlastní technologie pro integraci GPU do CPU. Z venku AMD Všechno
je to jasné: mají obrovské zkušenosti získané spolu s ATI.

Vytvoření „mikrozařízení“, Fusion, se bude skládat z několika jader
kódové označení Buldozer a video čip RV710 (Kong). Jejich vztah bude
prováděné prostřednictvím vylepšeného autobusu HyperTransport. Záleží na
počet jader a jejich frekvenční charakteristiky AMD plánuje vytvořit celou cenu
hierarchie "kamenů". Plánuje se také výroba procesorů pro notebooky (Falcon),
a pro multimediální gadgety (Bobcat). Navíc je to aplikace technologie
v přenosných zařízeních bude pro Kanaďany počáteční výzvou. S vývojem
paralelního počítání by mělo být používání takových „kamenů“ velmi populární.

Intel trochu pozadu v čase se svým Larrabee. produkty AMD,
pokud se nic nestane, objeví se na pultech obchodů koncem roku 2009 – brzy
2010. A rozhodnutí nepřítele vyjde najevo až za téměř dva
roku.

Larrabee bude mít velké množství (čti: stovky) jader. Na začátku
Nebudou chybět ani produkty určené pro 8 – 64 jader. Jsou velmi podobné Pentiu, ale
dost silně přepracován. Každé jádro má 256 kB L2 cache
(jeho velikost se časem zvětší). Vztah bude realizován prostřednictvím
1024bitová obousměrná kruhová sběrnice. Intel říká, že jejich „dítě“ bude
perfektně fungují s DirectX a Open GL API (pro vývojáře Apple), takže ne
není nutný žádný softwarový zásah.

Proč jsem ti to všechno řekl? Je zřejmé, že Larrabee a Fusion nevytlačí
běžné, stacionární procesory z trhu, stejně jako nebudou vytlačeny z trhu
grafické karty. Pro hráče a nadšence extrémních sportů zůstane konečný sen i nadále
vícejádrový CPU a tandem několika špičkových VGA. Ale co dokonce
procesorové společnosti přecházejí na paralelní výpočty založené na principech
podobně jako GPGPU, říká hodně. Zejména o tom, co takové
technologie jako CUDA má právo na existenci a s největší pravděpodobností bude
velmi populární.

Krátké shrnutí

Paralelní počítání pomocí grafické karty je jen dobrý nástroj
v rukou pracovitého programátora. Sotva procesory v čele s Moorovým zákonem
přijde konec. Společnosti NVIDIA k tomu je ještě dlouhá cesta
propagovat své API pro masy (totéž lze říci o duchovním dítěti ATI/AMD).
Jaké to bude, ukáže budoucnost. Takže CUDA se vrátí :).

P.S. Doporučuji návštěvu začínajícím programátorům a zájemcům
tyto „virtuální provozovny“:

Oficiální stránky a webové stránky NVIDIA
GPGPU.com. Všechno
uvedené informace jsou zapnuté anglický jazyk, ale alespoň děkuji, že ne
čínština Tak hurá do toho! Doufám, že vám autorka alespoň trochu pomohla
vzrušující cesty do průzkumu CUDA!