- procesor a GPU vždy asynchronní
- grid = celé gpu
- nekoherentní / nekonzistentní systém
- atomické operace jsou velice drahé
- architektura -> lepší hw škálovatelnost do budoucnosti (možnost volně přidávat více SM)
- in-order (v rámci 1 vlákna)
- SIMT architektura
- Vysoce úrovňové jazyky
- openACC - více high level, jinak obdobné openMP
- openMP a openACC - neumí zámky = neuděláme třeba největší prvek v poli
- Thrust - umí částečně STL podporu, např
std::for_each
- nemá přímou kontrolu
- pro aplikace, které nejsou moc složité (žádné šílené cachování, šílený přístup do paměti, …)
- CUDA??
- Nízko úrovňové jazyky
- CUDA
- OpenCL - obdobné CUDě, ale obecnější, musíme si postavit tu architekturu
- podpora NVIDII ustrnula asi před 10 lety
- HIP - v době kompilace si vybereme backend
- copmuta capability
- 6.0, 7.0 = server cards
- 6.1, 6.2, 7.5 = desktop cards
- Dobré si uložit query do nějakého logu - na staré/nepodporované kartě to nebude
nvidia-smi - checknout cuda verzi- integrated GPU = má sdílenou pamět s cpu, memcpy nedává smysl
- zavolání kernelu je pouze vložení do fronty
-
- spuštění trvá déle, kód je nakopírován do mezi kódu a karty si pak zvlášť dokompiluje pro sebe
__host__ __device__ před funkcí = funkce pro oboje cpu i gpu!__shared__ = alokace na cuda blocku (parallel shared cache)
- 1 block se vykonává vždy na 1 SM procesoru!!
- ale platí, že 1 SM procesor může vykonávat více blocků (podle počtu zdrojů)
- nedává smysl mít blocky po velikostech jiné, než dělitelné 32 (mít 32x8 radši než 16x16, bralo by to warpy po 2 řádcích)
- když nevíme, typicky 256 je dobrá startovací hodnota
- jde nám o ““maximum occupancy”“ - maximální
- omezovat zdroje na 1 block tak, aby bylo co největší využití zdrojů na 1 block
- slide Maximální využití SM