Pokročilé architektury počítačů

Pokročilé architektury počítačů Tutoriál 3

CUDA - GPU

Martin Milata

Výpočetní model CUDA ●

Organizace kódu ●

●

Sériově organizovaný kód určený pro CPU Paralelní kód prováděný na GPU –

●

GPU kernel ● ●

●

●

Označuje se jako kernel

provádí se na tzv. grid struktuře Grid se skládá z bloků (1D nebo 2D struktura) Blok je skupina až 512 vláken organizovaných do 1, 2 nebo 3 rozměrného pole

Hierarchizované uspořádání paralelismů

Paměťový model CUDA ●

Definuje různé adresní prostory pro přístup v rámci GPU a komunikaci s CPU ●

●

●

Globální paměť, paměť konstant a textur jsou přístupné GPU a CPU Registry, lokální a sdílená paměť přístupné GPU

Latence paměti ●

●

„Rychlá“ paměť s nízkou latencí – registry a sdílená paměť Frame buffer (DRAM paměť) – lokální paměť, paměť konstant a textur

GT200 ●

Více-jádrový čip s dvou vrstvou hierarchií ●

●

●

10x Thread Processing Cluster (TPC) na vyšší úrovni 3x Streaming Multiprocessor (SM) nebo Thread Processor Array (TPA) na nižší úrovni v rámci každého TPC

Hierarchie definovaná realizací přístupu do paměti ●

SM v rámci TPC sdílí přístup k hierarchií cache pamětí v rámci přístupu do paměti textur

Bližší pohled na výpočetní architekturu GT200 ●

Streaming Multiprocessor (SM) ●

SM Controller

●

Instrukční cache

●

Warp buffer

●

Tabulka skóre

●

8x Streaming Processor (SP) –

●

„Funkční jednotka“

2x speciální funkční jednotka – –

souvisí s grafickými výpočty lze je použít jako násobičky

●

1x 64bit funkční jednotka

●

16k 32bit registrů

●

16kB sdílené paměti

Třetí generace SM ●

32x CUDA Core ●

●

●

Konfiguračně rozdělena mezi L1 cache a sdílenou paměť (16KB x 48KB pro sdílenou paměť nebo cache)

16x L/S jednotka ●

●

ALU a FPU single i double precision

64KB Sdílené paměti ●

●

provádí celočíselné operace i operace v plovoucí řádové čárce

Jednotná cesta přístupu do paměti (dříve separována – paměť textur čtení a výstupní pixel zápis)

4x SFU ●

Speciální operace sin, cos, exp, rcp Obrázek převzat z: Whitepaper NVIDIA's Next Generation CUDA Compute Architecture: Fermi

Třetí generace SM plánování instrukcí

Obrázek převzat z: nvidia.com

Máte CUDA enabled zařízení?

Převzato z: http://www.nvidia.com/object/cuda_gpus.html Seznam není aktuální (verze z roku 2009)

Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
●

Stanovení velikosti vektoru ●

●

Počet prvků vektoru (10)

int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Nastavení počtu vláken bloku ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

square_array <<>>(a_d,N);

Maximální počet vláken v bloku (4)

cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);

Důsledek je popis velikosti gridu ●

Počet prvků / Počet vláken bloku }

for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Definice CUDA kernelu ●

●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

Realizuje paralelní výpočet druhých mocnin prvků vektoru na GPU

int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); square_array <<>>(a_d,N);

float *a – ukazatel do paměti GPU zařízení

cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);

int N – skutečný počet prvků

for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Definice CUDA kernelu ●

●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

Realizuje paralelní výpočet druhých mocnin prvků vektoru na GPU

int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); square_array <<>>(a_d,N);

float *a – ukazatel do paměti GPU zařízení

cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);

int N – skutečný počet prvků

for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Alokace a inicializace vektoru v hlavní paměti („paměť CPU“) ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

Vektor a_h velikosti N float hodnot je k dispozici v paměti CPU

int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

size jako celková velikost pole pro uložení N float hodnot bude použita i při následné alokaci GPU paměti

square_array <<>>(a_d,N); cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Alokace a kopírování hodnot vektoru do paměti GPU ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice);

Vektor a_d velikosti N float hodnot je alokován v paměti GPU

int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Při kopírování pamětí je specifikován a_d cíl a a_h zdroj přenosu, size velikost kopírované oblasti a cudaMemcpyHostToDevice směr přenosu

square_array <<>>(a_d,N); cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Výpočet počtu bloků vláken, do kterých bude provádění rozprostřeno ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Proměnná n_blocks definuje počet bloků v rámci gridu v 1D uspořádání

square_array <<>>(a_d,N);

Případný nenulový zbytek celočíselného podílu počtu prvků a počtu vláken v bloku vynutí alokaci bloku navíc (nebude plně využit)

cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost); for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Volání provádění CUDA kernelu ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Za voláním funkce kernelu následují parametry direktivy volání n_blocks počet bloků a block_size jejich velikost (počet vláken v bloku)

square_array <<>>(a_d,N); cudaMemcpy(a_h, a_d, sizeof(float)*N, cudaMemcpyDeviceToHost);

Vektor a_d je předán jako ukazatel do paměti GPU, následuje konstantní N

for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Kopírování hodnoty vypočtem modifikovaného vektoru zpět do paměťového prostoru CPU ●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); int n_blocks = N/block_size + (N%block_size == 0 ? 0:1);

Při kopírování pamětí je opět specifikován a_h cíl a a_d zdroj přenosu, size velikost kopírované oblasti a cudaMemcpyDeviceToHost směr přenosu

square_array <<>>(a_d,N); cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost); for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Zobrazení výsledku

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); square_array <<>>(a_d,N); cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost); for (int i=0; i
Příklad: Druhá mocnina prvků vektoru

const int N = 10; const int blocksize = 4;

__global__ void square_array(float *a,int N) { int idx=blockIdx.x*blockDim.x+threadIdx.x; if (idx
●

int main(void) { float *a_h; const size_t size = N*sizeof(float); a_h = (float *)malloc(size); for (int i=0; i
Úklid dynamicky alokovaných pamětí před ukončením provádění programu ●

●

cudaMemcpy(a_d, a_h, size, cudaMemcpyHostToDevice); int n_blocks = N/block_size + (N%block_size == 0 ? 0:1); square_array <<>>(a_d,N);

Uvolnění alokované paměti v prostoru GPU pomocí cudaFree(a_d)

cudaMemcpy(a_h, a_d, size, cudaMemcpyDeviceToHost); for (int i=0; i
Uvolnění paměti v prostoru CPU pomocí free(a_h) }

Vykonávání vláken ●

●

Struktura GPU ●

Obsahuje N multiprocesorů

●

Každý multiprocesor obsahuje M skalárních procesorů

Každý multiprocesor zpracovává skupiny bloků vláken ●

●

Bloky jsou rozděleny do skupin vláken tzv. warp ●

Warp je prováděn paralelně

●

V současné době obsahuje 32 vláken

●

●

Blok vláken může běžet jen na jednom multiprocesoru

Vlákna jsou ve warp řazena pokud možno se sekvenčně vzrůstajícím threadID

Plánovač přepíná provádění mezi warp instrukcemi

CUDA výpočetní server ●

●

Server Host:

tesla.cs.vsb.cz

Uživatel:

pap_cuda

Heslo:

cuda_test

Přihlášení ssh -X [email protected]

●

Kopírování zdrojového kódu na serveru mkdir ./ cp ./SRC/* .//

●

Kompilace CPU gcc -lrt <soubor>.c

GPU nvcc <soubor>.cu

CUDA práce s kódem ●

Editace kódu gedit ~//<soubor> &

●

Úkol ●

Získejte doby provádění kódu pro vyplnění následující tabulky CPU Inicializace (příprava matic) kopírování dat (přesun mezi pamětmi CPU a GPU) výpočet kopírování dat (přesun mezi pamětmi GPU a CPU

GPU

Literatura ●

D. Kanter: NVIDIA's GT200: Inside a Papallel Processor

●

NVIDIA CUDA C Programming Guide

●

Johan Seland: CUDA Programming

●

●

●

Paul H. J. Kelly, Advanced Computer Architecture Lecture notes 332 P. N. Glaskowsky: NVIDIA’s Fermi: The First Complete GPU Computing Architecture Internetové zdroje: ●

http://www.nvidia.com/

●

http://gpgpu.org/