CUDA

Téma prezentace: CUDA
Typ souboru: prezentace PDF
Přidal(a): Dabling
 

 
Popis materiálu:
Tato prezentace se zabývá základy technologie CUDA. Není v PPTX, ale PDF, takže spíše jako inspirace nebo učební materiál.
 
Obsah:
Obecné Informace Programový Model Compute Unified Device Architecture Brno, 2013 CUDA Obecné Informace Programový Model Vysvetlení pojm ˚u ˇ Pojem/zkratka význam CPU procesor GPU grafický procesor MS multiprocesor – procesory, ze kterých se skládá GPU kernel základní funkce provádející se na GPU ˇ VRAM/DRAM základní pamet’ grafické karty ˇ FLOPs pocet operací s desetinnou ˇ cárkou za sekundu ˇ host cip, na kterém je primárn ˇ e spušt ˇ ena aplikace (CPU) ˇ device zaˇrízení, na kterém beží CUDA (graf. karta) ˇ CUDA Obecné Informace Programový Model Motivace Základní architektura Motivace zkratka CUDA cipy spole ˇ cnosti nVidia ˇ technologie od roku 2006; nejstarší technologie pro paralelní výpocty na GPU ˇ dnes je podporována témeˇˇr v každé grafické karte s ˇ cipem ˇ spolecnosti nVidia ˇ multiplatformní technologie (Linux, Mac OS, Windows) podpora jazyk ˚u C/Cpp, FORTRAN, Python, Java, Open CL, DirectX Compute CUDA Obecné Informace Programový Model Motivace Základní architektura Motivace Proc grafické karty? ˇ CUDA Obecné Informace Programový Model Motivace Základní architektura Motivace Proc grafické karty? ˇ CUDA Obecné Informace Programový Model Motivace Základní architektura Motivace Proc grafické karty? ˇ renderování vektorové grafiky v reálném case ˇ ⇒ vysoké nároky na výkon architektura GPU se podstatne liší od CPU ˇ m ˚užeme této architektury využít i pro obecné výpocty ˇ Nekteré typy výpo ˇ ct ˚u nelze na GPU provád ˇ et efektivn ˇ eˇ využití: hromadné zpracování dat paralelní výpocty ˇ bruteforcing 🙂 CUDA Obecné Informace Programový Model Motivace Základní architektura Instalace http://docs.nvidia.com/cuda/ již je 5. verze CUDA instalacní balík obsahuje vše pot ˇ ˇrebné (nižší verze obsahovaly více r ˚uzných soubor ˚u) je nutné správne nalinkovat CUDA knihovny ˇ CUDA Obecné Informace Programový Model Motivace Základní architektura Hardwarová Architektura GPU je rozdeleno na tzv. multiprocesory (dnes Streaming ˇ Multiprocessor/CUDA Cores) každý MS je samostatná výpocetní jednotka ˇ 8 – 192 jader (liší se podle stáˇrí kary) 32bit registry (8K – 64K) sdílenou pamet’ (16KB – 48KB) ˇ konstantní pamet’ (64KB) ˇ texturovací pamet’ ˇ GPU komunikuje s CPU pomocí VRAM CUDA Obecné Informace Programový Model Motivace Základní architektura Hardwarová Architektura Základní schéma multiprocesoru CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Základní Programové Pojmy kernel – funkce definována klícovým slovem ˇ __global__ tato funkce je spuštena ˇ N-krát zdrojový kód již zpracovává N vláken každé vlákno má sv ˚uj index (threadIdx) tˇríprvkový vektor[x, y, z] – promenná typu dim3 (m ˚uže být ˇ nahrazena celocíselnou prom ˇ ennou) ˇ CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Volání Kernelu __global__ void Soucet(int ∗A, int ∗B, int ∗C) { int idx = threadIdx.x; \\ index vlakna C[idx] = A[idx] + B[idx]; } int main() { … \\ soucet N prvku matic A a B Soucet<<<1; N>>>(A_device, B_device, C_device); … } CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Volání Kernelu syntaxe: <<<1, N>>> – „Execution Configuration“ první argument, pocet blok ˚u v m ˇ ˇrížce druhý argument, pocet vláken v bloku ˇ obe dv ˇ e prom ˇ enné nabývají t ˇ ˇrí hodnot – pro každou dimenzi jinou velikost pˇri volání kernelu urcujeme topologii kernelu ˇ implicitne je každá hodnota sou ˇ ˇradnic blok ˚u i mˇrížky nastavena na hodnotu 1 <<<200, 256>>> – 200 blok ˚u, každý o 256 vláknech CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Volání Kernelu __global__ void Soucet(int A[N][M], int B[N][M], int C[N][M]) { int idx = threadIdx.x; int idy = threadIdx.y; C[idx][idy] = A[idx][idy] + B[idx][idy]; } int main() { … dim3 vlakna(N, M); Soucet<<<1; vlakna>>>(A_device, B_device, C_device); … } CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Užití Indexu Vláken index vlákna platí pouze v rámci bloku každý blok má vlastní ID v mˇrížce blockIdx[x, y, z] k vypocítání indexu pot ˇ ˇrebujeme velikost bloku blockDim[x, y, z] pokud chceme spustit na jedno pole více blok ˚u s vlákny, každé vlákno si musí spocítat index v poli, se kterým bude ˇ pracovat idx = blockIdx.x * blockDim.x + threadIdx.x; CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Volání Kernelu __global__ void Soucet(int A[N][M], int B[N][M], int C[N][M]) { int idx = blockIdx.x ∗ blockDim.x + threadIdx.x; int idy = blockIdx.y ∗ blockDim.y + threadIdx.y; C[idx][idy] = A[idx][idy] + B[idx][idy]; } int main() { … dim3 vlaknaNaBlok(16, 16); //256 dim3 pocetBloku(N / vlaknaNaBlok.x, M / vlaknaNaBlok.y); Soucet<<<pocetbloku, vlaknanablok=““>>>(A_device, B_device, C_device); … } CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Základy Práce s Pametí ˇ rozlišujeme RAM a VRAM rozšíˇrená syntaxe standardní alokace pameti GPU ˇ cudaMalloc((void**) &p_device, size); p_device musí být pointer data je nutné na GPU pˇrekopírovat cudaMemcpy(p_device, p_host, size, cudaMemcpyHostToDevice); cudaMemcpy(p_host, p_device, size, cudaMemcpyDeviceToHost); uvolnení pam ˇ eti ˇ cudaFree(p_device); hodnota – cudaError_t CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Hello CUDA! prográmek 🙂 CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Bloky, Mˇrížky, Vlákna r ˚uzná GPU mají r ˚uzný pocet MP ˇ není nutno rekompilovat program kv ˚uli r ˚uzným GPU sdílená pamet’, synchronizace vláken v rámci jednoho ˇ bloku CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Pameti Podrobn ˇ eji ˇ Typ pameti ˇ Pˇrístup Umístení ˇ Operace Keš. Registry Jedno vlákno Cip ˇ ctení/zápis ˇ Ne Lokální Jedno vlákno DRAM ctení/zápis ˇ Ne Sdílená Vlákna v rámci bloku Cip ˇ ctení/zápis ˇ – Globální Všechna vlákna + CPU DRAM ctení/zápis ˇ Ne Texturovací Všechna vlákna + CPU DRAM ctení/zápis ˇ Ano Konstantní Všechna vlákna + CPU DRAM GPU jen ctení ˇ Ano CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Sdílená Pamet’ ˇ globální pamet’ je p ˇ ˇríliš pomalá, sdílená pamet’ je ˇ mnohonásobne rychlejší ˇ vlákna v jednom bloku mohou spolupracovat, jsou propojena sdílenou pametí ˇ žádná jiná vlákna nemají do sdílené pameti p ˇ ˇrístup po zániku bloku zaniká i alokovaná sdílená pamet’ ˇ užití: pole s cast ˇ ejším využitím v rámci jednoho bloku ˇ KONFLIKTY BANK – nacítání hodnot, které ješt ˇ e neexistují ˇ CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Alokace Sdílené Pameti ˇ v Execution Configuration – tˇretí argument __global__ kernel() { extern __shared__ int SHARED_MEMORY[]; } int main() { … kernel<<<bloky, vlakna,=““ velikost_sdilene_pameti_v_bytech=““>>>(); … } CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Alokace Sdílené Pameti ˇ v kernelu __global__ kernel() { __shared__ int SHARED_MEMORY[velikost_v_bytech]; } int main() { … kernel<<<bloky, vlakna=““>>>(); … } CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Sdružený Pˇrístup Do Pameti ˇ Coalesced Memory Access globální pamet’ je veliká, ale pomalá a nekešovaná ˇ je zde technika pro využívání paralelních cinností vláken ˇ každá instrukce je zpracovávána ve warpech; 32 vláken vlákna ve warpu nesmí divergovat! ⇒ zvýšení rychlosti delení na Half-Warpy (16 vláken) ˇ pro sdružený pˇrístup platí tˇri pravidla: velikost promenných
musí být 4, 8, 16 byt ˚u ˇ vlákna k promenným p ˇ ˇristupují sekvencnˇ e (N-té vlákno ˇ pˇristupuje k N-tému prvku) promenné musí být ve stejném pam ˇ et’ovém segmentu, ˇ adresa prvního prvku je zarovnána k 16ti násobku velikosti promenné ˇ CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Sdružený Pˇrístup Do Pameti ˇ Colesced Memory Access výrazné snížení latencí pameti ˇ promenné je t ˇ ˇreba upravit podle zmínených velikostí ˇ CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Pole v CUDA optimální jsou jednorozmerná pole ˇ 2D a 3D pole nelze v CUDA použít cudaMallocPitch() a cudaMalloc3D() (cudaMemcpy2D() a cudaMemcpy3D()) funkce nutné kv ˚uli nacítání blok ˚u dat ˇ malloc() vs cudaMallocHost(); CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Proudy Streams chceme více paralelizovat! chceme zamestnávat CPU i GPU zárove ˇ nˇ proud – posloupnost pˇríkaz ˚u, které se provádejí v jedné ˇ fronteˇ m ˚užeme vytvoˇrit více front cíl: CPU pˇred-zpracovává data a spouští kernely FSB kopíruje data z/na GPU GPU zpracovává data CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Proudy Streams rozdelit problém na menší ˇ cásti ˇ každou cást provád ˇ et postupn ˇ eˇ viz nVidia prezentace a pˇríklad … CUDA Obecné Informace Programový Model Volání Kernelu Pamet’ ˇ Proudy Synchronizace __syncthreads() cudaThreadSynchronize() cudaDeviceSynchronize() cudaStreamSynchronize() cudaStreamWaitEvent() cudaStraeamQuery() cudaEventSynchronize() cudaEventQuery() CUDA