GPU hardware Pamětoua hierarchie Synchronizace Násobeni matic GPU hardware a paralelismus Jiří Filipovič podzim 2009 tA-Mmmiumnum GPU hardware Pamětoua hierarchie Synchronizace Násobeni matic Alternativy k CUDA CUDA je (a zřejmě i bude) pouze pro GPU nVidia. RagSBI □ SP - = -^ ^o^o GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Alternativy k CUDA CUDA je (a zřejmě i bude) pouze pro GPU nVidia. OpenCL • standard, pro různé druhy akcelerátorů (nezávislý na výrobci HW i operačním systému) • silně inspirováno CUDA, velmi snadný přechod • z principu nemůže reagovat tak rychle na vývoj HW s ■O Q-C^ GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Alternativy k CUDA CUDA je (a zřejmě i bude) pouze pro GPU nVidia. OpenCL • standard, pro různé druhy akcelerátorů (nezávislý na výrobci HW i operačním systému) • silně inspirováno CUDA, velmi snadný přechod • z principu nemůže reagovat tak rychle na vývoj HW DirectX compute • GPU různých výrobců, OS pouze jednoho s ■O Q-C^ GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Alternativy k CUDA CUDA je (a zřejmě i bude) pouze pro GPU nVidia. OpenCL • standard, pro různé druhy akcelerátorů (nezávislý na HW i operačním systému) • silně inspirováno CUDA, velmi snadný přechod • z principu nemůže reagovat tak rychle na vývoj HW DirectX compute • GPU různých výrobců, OS pouze jednoho Brook(+) • n ad platform ní, + jen pro AM D /AT I • pouze streamy s GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Proč se učíme o CUDA? Dnes platí • největší množství aplikací • největší množství knihoven • největší počet publikací • snadné k naučení • podobnost s OpenCL umožňuje snadný přechod • i v n ad platform ním jazyce je dobré přizpůsobovat se jednotlivým platformám s ■O Q-C^ GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Rozdíly v jednotlivých CUDA-GPU Nové generace přinášejí vyšší výkon a výpočetní schopnosti. • výpočetní schopnosti compute capability představují bohatost instrukční sady GPU a množství zdrojů (registry, současně běžící thready aj.) o výkon roste se schopností umístit na jedno GPU více jader V rámci generace se významně liší výkon. • kvůli nabídce levnějších variant • díky pozdějším změnám výrobního procesu • kvůli spotřebě u mobilních GPU GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Dnes dostupné GPU Dnes dostupné GPU • compute capability 1.0 - 1.3 • s rozdíly se budeme postupně seznamovat • 1 - 30 multi procesorů (19.2 - 1 062.7 GFLOPs) • frekvence 800 MHz - 1.836GHz • šířka a rychlost datové sběrnice (64bit - 512bit, 6.4 159 GB/s) S •f) <\(y MJII.UJ.II.IJ1I.UJ.IIJIU..IW GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Dostupná řešení Grafické karty GeForce • mainstreamové řešení pro hráče • levné, široce rozšířené, široké spektrum výkonu • nevýhoda - omezená paměť (do 1 GB na GPU) Profesionální karty Quadro • z hlediska CUDA stejné jako GeForce • paměť až 4 GB na GPU • násobně dražší Tesla • řešení speciálně pro výpočty v CUDA • jedno GPU na generaci (základní varianta), vždy velká paměť • k dispozici jako karty do PCI-E, nebo jako samostatné multi-GPU stroje • taktéž drahé, vhodné pro výpočetní centra či osobní superpočítače •f) <\(y Paralelismus GPU Paralelní algoritmy je nutno navrhovat vzhledem k paralelismu, které poskytuje HW • v případě GPU se jedná o pole SIMT m u Iti procesorů pracující nad společnou pamětí Dekompozice pro GPU • hrubozrnne rozdělení problému na části nevyžadující intenzivní komunikaci/synchronizaci • jemnozrnné rozdělení blízké vektorizaci (SIMT je ale více flexibilní) GPU hardware Paralelismus Paměťová hierarcl Synchronizace Násobení matic Hierarchie vláken Grid Block (0,0) Block (1,0) Block (2,0) Block (0,1)- Block (1,1) v Block (2,1) Three Block (1,1) zl Thread (0, 0) Thread (1, 0) Thread (2, 0) Thread (3, 0) Thread (0,1) Thread (1,1) Thread (2,1) Thread (3,1) Thread (1,2) Thread (3, 2) □ &1 - Jiří Filipovič GPU hardware a paralelismus GPU hardware Pamětova hierarchie Synchronizace Násobeni matic SIMT Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci iMmumiumnum •f)<\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic SIMT Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít GPU hardware Pamětova hierarchie Synchronizace Násobeni matic SIMT Multiprocesor má jen jednu jednotku pro spouštějící instrukce • všech 8 SP musí provádět stejnou instrukci • nová instrukce je spuštěna každé 4 cykly • 32 thredů (tzv. warp) musí provádět stejnou instrukci A co větvení kódu? • pokud část threadů ve warpu provádí jinou instrukci, běh se serializuje • to snižuje výkon, snažíme se divergenci v rámci warpu předejít Multiprocesor je tedy MIMD (Multiple-Instruction Multiple-Thread) z programátorského hlediska a SIMT (Single-Instruction Multiple-Thread) z výkonového. Paměťová hierarchie Synchronizace Násobení matic GPU architektura □ g - = š "O^O Jiří Filipovič GPU hardware a paralelismus GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Vlastnosti threadů Oproti CPU threadům jsou GPU thready velmi lehké (lightweight). • jejich běh může být velmi krátký (desítky instrukcí) • může (mělo by) jich být velmi mnoho • nemohou využívat velké množství prostředků Thready jsou seskupeny v blocích • ty jsou spouštěny na jednotlivých m u Iti procesorech • dostatečný počet bloků je důležitý pro škalovatelnost Počet threadů a thread bloků na multiprocesor je omezen. s ■O Q-C^ GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci s ■O Q-C^ MJII.UJ.II.IJ1I.UJ.IIJIU..IW GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache s GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overhead u GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Maskování latence pamětí Paměti mají latence • globální paměť má vysokou latenci (stovky cyklů) • registry a sdílená paměť mají read-after-write latenci Maskování latence paměti je odlišné, než u CPU • žádné provádění instrukcí mimo pořadí • u většiny pamětí žádná cache Když nějaký warp čeká na data z paměti, je možné spustit jiný • umožní maskovat latenci paměti • vyžaduje spuštění řádově více vláken, než má GPU jader • plánování spuštění a přepínání threadů je realizováno přímo v HW bez overhead u Obdobná situace je v případě synchronizace. Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) Registry • nejrychlejší paměť, přímo využitelná v instrukcích • lokální proměnné v kernelu i proměnné nutné pro mezivýsledky jsou automaticky v registrech • pokud je dostatek registrů • pokud dokáže kompilátor určit statickou indexaci polí • mají životnost threadu (warpu) Lokální pamět • co se nevleze do registrů, jde do lokální paměti • ta je fyzicky uložena v DRAM, je tudíŠ pomalá a má dlouhou latenci • má životnost threadu (warpu) Sdílená paměť • rychlá jako registry • nedojde-li ke konfliktům paměťových bank • instrukce umí využít jen jeden operand ve sdílené paměti (jinak je třeba explicitní load/store) • v C for CUDA deklarujeme pomocí shared— • proměnná ve sdílené paměti může mít dynamickou velikost (určenou při startu), pokud je deklarována jako extern bez udání velikosti pole • má životnost bloku Deklarace statické sdílené paměti __shared__ float myArray[128]; Dynamická alokace extern __shared__ char myArray []; float *arrayl = (float*)myArray; int *array2 = (int*)&array1[128]; short *array3 = (short*)&array2[256]; Vytvoří pole arrayl typu float velikosti 128, pole array2 typu int velikosti 256 a pole array3 plovoucí velikosti. Celkovou velikost je nutné specifikovat při spouštění kernelu. myKernel<«grid , block, n>>>(); Globální paměť • řádově nižší přenosová rychlost než u sdílené paměti • latence ve stovkách GPU cyklů • pro dosažení optimálního výkonu je třeba paměť adresovat zarovnaně • má životnost aplikace Lze dynamicky alokovat pomocí cudaMalloc, či staticky pomocí deklarace —device— Paměť konstant • pouze pro čtení • cacheována • cache-hit poskytuje rychlost jako registry (za dodržení určitých podmínek), cache-miss rychlý jako globální pamět • omezená velikost (64 KB u v současné dobe dostupných GPU) • má životnost aplikace V deklaraci používáme —constant—, ke kopírování dat do paměti konstant slouží funkce cudaError_t cudaMemcpyToSymbol(const char »symbol, const void *src, size_t count, size_t offset, enum cudaMemcpyKind kind) Kopíruje data ze systémové (cudaMemcpyHostToDevice) nebo globální paměti (čudaMemcpyDeviceToDevice) z src do symbol. Kopírovaný blok má velikost count bytů, do paměti symbol kopírujeme s posuvem offset. Paměť textur • cacheovaná, 2D prostorová lokalita • pouze pro čtení (omezení kvůli cache-koherenci) • dlouhá latence • více adresovacích módů • možnost normalizace dimenzí do [0.1] • při adresaci mimo rozsah ořezávání či přetékání koordinát • možnost filtrace dat • lineární interpolace nebo nejbližší hodnota • funkcionalita je ,,zdarma" (implementováno v HW) Více detailů viz CUDA Programming Guide. Systémová paměť RAM • s GPU spojena přes PCI-E • virtuální adresace komplikuje přenosy mezi CPU (host) a GPU (device) pamětí • je možno alokovat tzv. page-locked oblast paměti • může redukovat celkový výkon systému • omezená velikost • data jsou po PCI-E přenášena rychleji • umožňuje paralelní běh hernelu a kopírování dat • umožňuje namapovat adresovací prostor host paměti na device • umožňuje write-combining přístup (data nejsou cacheována CPU) Namísto mallocQ použijeme pro alokaci cudaMallocHost(), pro uvolnění cudaFreeHostQ • flag cudaHostAllocPortable zajistí vlastnosti page-locked paměti pro všechny (CPU) thready • flag cudaHostAllocWriteCombined vypne cacheování alokované paměti CPU • flag cudaHostAllocMapped nastaví mapování host paměti v device paměťovém prostoru Mapovaná paměť • totéž místo má rozdílnou adresu pro device a host kód • adresu pro device získáme pomocí cud a Host GetDevicePoin ter() • před voláním ostatních CUDA API funkcí je zapotřebí zavolat cudaSetDeviceFlagsQ s flagem cudaDeviceMapHost Asynchronní přenosy • API funkce mají suffix Async • může se překrývat přenos dat a výpočet na CPU i přenos dat a výpočet na GPU (podrobněji až probereme streamy) Necacheovaná paměť • pomalé čtení z host kódu • rychlejší přístup z device paměti Synchronizace v rámci bloku • nativní bariérová synchronizace • musí do ní vstoupit všechny thready (pozor na podmínky!) • pouze jedna instrukce, velmi rychlá, pokud neredukuje paralelismus • v C for CUDA volání __syncthreads() • komunikace přes sdílenou paměť • thready si přes ni mohou vyměňovat data • synchronizace atomickými operacemi, nebo bariérou s ■O Q-C* Atomické operace • provádí read-mod ify-write operace nad sdílenou nebo globální pamětí • žádná interference s ostatními thready • pouze pro celá 32-bitová či 64-bitová (pro compute capability > 1.2) čísla • nad globální pamětí u zařízení s compute capability > 1.1, nad sdílenou c.c. > 1.2 • aritmetické (Add, Sub, Exch, Min, Max, lne, Dec, CAS) a bitové (And, Or, Xor) operace Hlasování warpu Všechny thready v jednom warpu vyhodnocují podmínku a provedou její srovnání. Dostupné u zařízení s c.c. > 1.2. int __all(int predicate); Nabývá nenulové hodnoty tehdy a jen tehdy když je nenulový predikát pro všechny thready ve warpu. int __any(int predicate); Nebývá nenulové hodnoty tehdy a jen tehdy když alespoň jeden thread ve warpu vyhodnotí predikát jako nenulový. Synchronizace paměťových operací Sdílenou paměť obvykle využíváme ke komunikaci mezi thready a nebo jako cache pro data užívaná více thready. • thready využívají data uložené jinými thready • je třeba zajistit, abychom nečetli data, která ještě nejsou k dispozici • chceme-li počkat, až jsou data k dispozici, používáme syncthreadsQ Synchronizace paměťových operací Kompilátor může optimalizovat operace se sdílenou/globální pamětí (mezivýsledky mohou zůstat v registrech) a může měnit jejich pořadí, • chceme-li se ujistit, že jsou námi ukládaná data viditelná pro ostatní, používáme —threadfenceQ, popř. —thread fence-block() • deklarujeme-li proměnnou jako volatile, jsou veškeré přístupy k ní realizovány přes load/store do sdílené či globální paměti Synchronizace bloků Mezi bloky • globální paměť viditelná pro všechny bloky • slabá nativní podpora synchronizace • žádná globální bariéra • u novějších GPU atomicke operace nad globální pamětí • globální bariéru lze implementovat voláním kernelu (jiné řešení dosti trikové) • slabé možnosti globální synchronizace znesnadňují programování, ale umožňují velmi dobrou škalovatelnost GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Globální synchronizace přes atomické operace Problém součtu všech prvků vektoru • každý blok sečte prvky své části vektoru • globální bariéra • jeden blok sečte výsledky se všech bloků BOBBI S •f) <\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic _device__ unsigned int count = 0; .shared., bool isLastBlockDone; _global__ void sum(const float* array, unsigned int N, float* result) { float partialSum = calculatePartialSum(array, N); if (threadldx.x = 0) { result[blockldx.x] = partialSum; __threadfence(); unsigned int value = atomiclnc (&count , gridDim.x); isLastBlockDone = (value = (gridDim.x — 1)); } __syncthreads (); if (isLastBlockDone) { float totalSum = calculateTotalSum(result); if (threadldx.x == 0) { result[0] = totalSum; count = 0; } BOBBI a •f)<\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Násobení matic Chceme vynásobit matice A a B a výsledek uložit do C. Pro jednoduchost uvažujme matice velikosti n x n. CiJ = ELl Ai,k ■ Bk,j Zápis v jazyce C for (int i = 0; i < n; i++) for (int j = 0; j < n; j++){ C[i*n + j] = 0.0; for (int k = 0; k < n; k++) C[i*n + j] += A[i*n + k] * B[k*n + j]; } RagSBI S •f)<\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Paralelizace for (int i = 0; i < n; i++) for (int j = 0; j < n; j++){ C[i*n + j] = 0.0; for (int k = 0; k < n; k++) C[i*n + j] += A[i*n + k] * B[k*n + j]; } Problém lze paralelizovat více způsoby • vybrat jeden z cyklů • vybrat dva z cyklů • paralelizovat všechny cykly 9BBBSH •f) <\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) s ■O Q-C^ MJII.UJ.II.IJ1I.UJ.IIJIU..IW GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálování se zdá dobrá, počet vláken roste kvadraticky vzhledm k n s ■O Q-C^ GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Paralelizace Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálování se zdá dobrá, počet vláken roste kvadraticky vzhledm k n Paralelizace přes vnitřní cyklus • nevhodná, nutná serializace zápisu do C! s ■O Q-C^ Paralelizace přes jeden cyklus • neškáluje dobře, nutno používat velké matice (nezapomínejme, pro dobré využití GPU potřebujeme tisíce threadů) Paralelizace přes dva cykly • z hlediska škálování se zdá dobrá, počet vláken roste kvadraticky vzhledm k n Paralelizace přes vnitřní cyklus • nevhodná, nutná serializace zápisu do C! Jako nejvhodnější se tedy jeví paralelizovat cykly jdoucí přes / a j. S výhodou můžeme využít možnosti uspořádání bloku a mřížky jako 2D pole. __global__ void mmul(float *A, float *B, float *C, int n){ int x = blockldx.x*blockDim.x + threadldx.x; int y = blockldx.y*blockDim.y + threadldx.y; float tmp = 0; for (int k = 0; k < n; k++) t mp += A[y * n+k ] * B[k * n+x] ; C[y*n + x] = tmp; } Povšimněte si nápadné podobnosti s matematickým zápisem -paralelní verze je intuitivnější, než sériová! □ s ~ - GPU hardware Výkon Pamětoua hierarchie Synchronizace Násobeni matic Jaký bude mít naše implementace výkon? □ ► « ĚP ► < -^ ► 4 -š ► MJII.UJ.II.IJ1I.UJ.IIJIU..IW GPU hardware Výkon Pamětova hierarchie Synchronizace Násobeni matic Jaký bude mít naše implementace výkon? Uvažujme kartu GeForce GTX 280 • pro problém násobení matic využitelných 622GFIops • propustnost paměti 142 GB/s s ■O Q-C^ MJII.UJ.II.IJ1I.UJ.IIJIU..IW Jaký bude mít naše implementace výkon? Uvažujme kartu GeForce GTX 280 • pro problém násobení matic využitelných 622GFIops • propustnost paměti 142 GB/s Flop-to-word ratio naší implementace • v jednom kroku cyklu přes k načítáme 2 floaty (jedno číslo z matice A, jedno z B) a provádíme dvě aritmetické operace • jedna aritmetická operace připadne na přenos jednoho floatu (4 bytů) • globální paměť má propustnost 35.5 miliardy floatu za sekundu, pokud jeden warp přenáší jeden float z jedné matice a 16 floatu z druhé, můžeme dosáhnout výkonu cca GPU hardware Pamětoua hierarchie Synchronizace Násobeni matic Co s tím? Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? BOBBI S •f) <\(y GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Co s tím? Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? Pro výpočet jednoho prvku C musíme načíst řádek z A a sloupec z B, které jsou v globální paměti. BOBBI S •f) <\(y Narazili jsme na výkon globální paměti. GPU poskytuje rychlejší paměti, můžeme je využít? Pro výpočet jednoho prvku C musíme načíst řádek z A a sloupec z B, které jsou v globální paměti. Je opravdu nutné dělat to pro každý prvek C zvlášť? • pro všechny prvky C ve stejném řádku načítáme stejný řádek A • pro všechny prvky C ve stejném sloupci načítáme stejný sloupec B • můžeme některá data načíst jednou z globální paměti do sdílené a následně je opakovaně číst z rychlejší sdílené paměti Budeme-li přistupovat k matici po blocích, můžeme amortizovat přenosy z globální paměti: • počítáme část matice C o velikosti a x a • iterativně načítáme bloky stejné velikosti z matic A a B do sdílené paměti • tyto bloky spolu násobíme a výsledek přičítáme do C • poměr aritmetických operací k přeneseným floatům je a-násobný Přirozené mapování na paralelismus v GPU • jednotlivé bloky threadů budou počítat bloky matice C • mají společnou sdílenou paměť • rychle se synchronizují • mezi bloky není třeba žádné synchronizace Jak velké bloky zvolit? • jsme omezeni velikostí sdílené paměti • jsme omezeni počtem threadů, který může běžet na GPU • necháme-li jeden thread počítat jeden prvek C, jeví se jako rozumná velikost bloku 16 x 16 • jedná se o násobek velikosti warpu • jeden blok bude mít únosných 256 threadů • jeden blok spotřebuje 2 KB sdílené paměti • paměť nebude zásadně omezovat výkon (16 • 25.5 = 568 GFIops, což je již poměrně blízko hodnotě 622 GFIops) Schéma algoritmu • každý blok threadů bude mít pole As a Bs ve sdílené paměti • iterativně se budou násobit bloky matic A a B, výsledek bude každý thread kumulovat v proměnné Csub • thready v bloku společně načtou bloky do As a Bs • každý thread vynásobí bloky v As a Bs pro jeho prvek výsledné matice v Csub • každý thread uloží jeho prvek matice do globální paměti C Pozor na synchronizaci • než začneme násobit jednotlivé bloky, musí být kompletně načteny • než začneme znovunačítat bloky, musí být dokončeno násobení s původními daty GPU hardware Pamětova hierarchie Synchronizace Násobeni matic Druhý kernel _global__ void mmul(float *A, float *B, float *C, int n){ int bx = blockldx.x; int by = blockldx.y; int tx = threadldx.x; int ty = threadldx.y; __shared__ float As[BLOCK.SIZE][BLOCK.SIZE]; __shared__ float Bs[BLOCK.SIZE][BLOCK.SIZE]; float Csub = O.Of; for (int b = 0; b < n/BL0CK_SIZE; b++){ As[ty][tx] = A[(ty + by*BL0CK_SIZE)*n + b*BL0CK_SIZE+tx]; Bs[ty][tx] = B[(ty + b*BL0CK_SIZE)*n + bx*BLOCK_SIZE+tx]; __syncthreads(); for (int k = 0; k < BLOCK.SIZE; k++) Csub += As[ty][k]*Bs[k][tx]; __syncthreads(); } C[(ty + by*BL0CK)*n + bx*BLOCK_SIZE+tx] Csub ; S ■O Q-C^ • teoretické omezení první verze kernelu je 66.8 GFIops, naměřený výkon 36.6 GFIops • teoretické omezení druhé verze kernelu je 568 GFIops, naměřený výkon 198 GFIops • jak se přiblížit maximálnímu výkonu karty? • je třeba znát podrobněji HW a jeho omezení a podle toho navrhovat algoritmy • látka na další přednášku :-) s ■O Q-C^