GPGPU. Jan Faigl. Gerstnerova Laboratoř pro inteligentní rozhodování a řízení České vysoké učení technické v Praze

GPGPU Jan Faigl Gerstnerova Laboratoˇr pro inteligentní rozhodování a ˇrízení ˇ Ceské vysoké uˇcení technické v Praze

8. cviˇcení

ˇ katedra kybernetiky, FEL, CVUT v Praze

X33PTE - Programovací techniky – GPGPU

1 / 14

Masivní paralelizmus v grafických kartách • Výpoˇcet (rendering) obrazu probíhající po jednotlivých

pixelech lze velmi snadno paralelizovat. • Dedikované grafické procesoru (GPU), vysoký stupenˇ

integrace srovnatelný s hlavními procesory. • Vysoký poˇcet paralelních procesoru. ˚ • Využití výpoˇcetního výkonu v jiných aplikacích: • Zpracování proudu˚ dat (SIMD instrukce - procesory). • GPGPU - General Purpose computation on GPU. http://www.gpgpu.org • OpenCL (Open Computing Language) - abstrakce nad rozhraními GPGPU. • CUDA - rozhraní pro grafické karty spoleˇcnosti NVIDIA. http://www.nvidia.com/object/cuda_home.html



2 / 14

Výkon procesoru˚ • Jaký je udávaný výkon procesoru? ˚ • Grafické (streamovací) procesory CSX700 Cell GeForce 8800 GTX Radeon HD 4670

96 GigaFLOPs 102 GigaFLOPs 518 GigaFLOPs 480 GigaFLOPs

katalogové špiˇckové hodnoty. • Hlavní procesory: Phenom X4 9950 (@2.6 GHz) Core 2 Duo E8600 (@3.3 GHz) Cure 2 Quad QX9650 (@3.3 GHz) Cure 2 Quad QX9650 (@3.3 GHz) Core i7 967 (@3.2 GHz)

21 GigaFLOPs 22 GigaFLOPs 35 GigaFLOPs 35 GigaFLOPs 42 GigaFLOPs

Test linpack 32-bit (zdroj: www.pctuning.cz). • Je udávaný výkon dosažitelný? (float vs double) • Co jiné metriky, napˇr. výkon / spotˇreba. (CSX700 typická spotˇreba 9W) ˇ katedra kybernetiky, FEL, CVUT v Praze


3 / 14







3 / 14







3 / 14







3 / 14







3 / 14







3 / 14

CUDA • NVIDIA Compute Unified Device Architecture. • Rozšíˇrení syntaxe C o pˇrístup k paralelním jednotkám

GPU. • Výpoˇcet (kernel) je prováden ˇ jednotkou GPU. • Kernel je prováden ˇ paralelneˇ na více jednotkách. • Host - hlavní procesor (proces). • Device - procesor GPU. • Data musejí být v pameti ˇ pˇrístupné GPU, pˇresun

ˇ → Device pamet’. ˇ Host pamet’ • Výsledek výpoˇctu je v pameti ˇ GPU, pˇresun

ˇ ← Device pamet’. ˇ Host pamet’ ˇ katedra kybernetiky, FEL, CVUT v Praze


4 / 14

CUDA - organizace výpoˇctu

• Výpoˇcet (kernel) je rozdelen ˇ na bloky. • Blok reprezentuje paralelní výpoˇcet cˇ ásti výsledku. Napˇr. cˇ ást souˇcinu dvou matic. • Blok se skládá z výpoˇcetních vláken. • V rámci bloku mohou být paralelní výpoˇcty

synchronizovány. • Bloky jsou organizovány do gridu. • Škálovatelnost je realizována rozdelením ˇ výpoˇctu do bloku. ˚

ˇ záleží Blok resp. bloky nemusejí být nutneˇ poˇcítány paralelne, na konkrétním HW a skuteˇcném poˇctu paralelní jednotek GPU.



5 / 14

ˇ CUDA - grid, bloky, vlákna a pˇrístup do pameti Host − CPU

Device − GPU Grid 1

Kernel 1

Grid

Blok

Blok

(0, 0)

(1, 0)

Blok

Blok

(0, 1)

(1, 1)

Blok (0, 0)

Grid 2 Kernel 2

Blok

Blok

(0, 0)

(1, 0)

(2, 0)

Blok Blok (1, Blok 1)

Blok

(0, 1)

(2, 1)

Vlákno

Vlákno

Vlákno

(0, 0)

(1, 0)

(2, 0)

Vlákno

Vlákno

Vlákno

(0, 1)

(1, 1)

(1, 2)


Sdílená paměť

Registry

Registry

Registry

Registry

Vlákno (0, 0)

Vlákno (1, 0)

Vlákno (0, 0)

Vlákno (1, 0)

Lokální paměť

Blok

(1, 1)

Blok (1, 0)

Sdílená paměť

Lokální paměť

Lokální paměť

Lokální paměť

Globální paměť Paměť konstant Paměť textur

• Rychlost pˇrístupu do pameti. ˇ • Kolize pˇrístup více vláken.


6 / 14

CUDA - pˇríklad - Násobení matic 1/8 • NVIDIA CUDA SDK - Version 2.0, matrixMul. • Jednoduché násobení matic: • C = A · B, • matice jsou rozmeru ˇ n × n, • kde n je násobek velikosti bloku. • Porovnání: • naivní implementace v C (3× for smyˇcka), • naivní implementace v C s transpozicí, • CUDA implementace. • Hardware: • CPU - Intel Core 2 Duo @ 3 GHz, 4 GB RAM, • GPU - NVIDIA G84 (GeForce 8600 GT), 512 MB RAM.



7 / 14

CUDA - pˇríklad - Násobení matic 2/8

Naivní implementace 1 2 3 4 5 6 7 8 9 10 11

void simple_multiply(const int n, const float *A, const float *B, float *C) { for (int i = 0; i < n; ++i) { for (int j = 0; j < n; ++j) { float prod = 0; for (int k = 0; k < n; ++k) { prod += A[i * n + k] * B[k * n + j]; } C[i * n + j] = prod; } } }



8 / 14

CUDA - pˇríklad - Násobení matic 3/8 Naivní implementace s transpozicí 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21

void simple_multiply_trans(const int n, const float const float *b, float *c) { float * bT = create_matrix(n); for (int i = 0; i < n; i++) { bT[i*n + i] = b[i*n + i]; for (int j = i + 1; j < n; j++) { bT[i*n + j] = b[j*n + i]; bT[j*n + i] = b[i*n + j]; } }

*a,

for (int i = 0; i < n; i++) { for (int j = 0; j < n; j++) { float tmp = 0; for (int k = 0; k < n; k++) { tmp += a[i*n + k] * bT[j*n + k]; } c[i*n + j] = tmp; } } free(bT); }



9 / 14

CUDA - pˇríklad - Násobení matic 4/8 CUDA - strategie výpoˇctu: B

• rozdelení ˇ matice na bloky, • každý blok vypoˇcte jednu

sub-matici Csub , A

C

vypoˇcte jeden element Csub .

C sub

BLOCK_SIZE

• každé vlákno v bloku

BLOCK_SIZE



10 / 14

CUDA - pˇríklad - Násobení matic 5/8 CUDA implementace - hlavní funkce 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23

void cuda_multiply(const int n, const float *hostA, const float *hostB, float *hostC) { const int size = n * n * sizeof(float); float *devA, *devB, *devC; cudaMalloc((void**)&devA, size); cudaMalloc((void**)&devB, size); cudaMalloc((void**)&devC, size); cudaMemcpy(devA, hostA, size, cudaMemcpyHostToDevice); cudaMemcpy(devB, hostB, size, cudaMemcpyHostToDevice); //BLOCK_SIZE == 16 dim3 threads(BLOCK_SIZE, BLOCK_SIZE); dim3 grid(n / threads.x, n /threads.y); //volání kernel funkce matrixMul matrixMul<<>>(n, devA, devB, devC); cudaMemcpy(hostC, devC, size, cudaMemcpyDeviceToHost); cudaFree(devA); cudaFree(devB); cudaFree(devC); }



11 / 14

CUDA - pˇríklad - Násobení matic 6/8 CUDA implementace - kernel funkce 1 2 3 4 5 6 7 8 9 10 11 12 13 14 15 16 17 18 19 20 21 22 23 24

__global__ void matrixMul(int n, float* A, float* B, float* C) { int bx = blockIdx.x; int by = blockIdx.y; int tx = threadIdx.x; int ty = threadIdx.y; int aBegin = n * BLOCK_SIZE * by; //zaˇ cátek sub-matice v bloku int aEnd = aBegin + n - 1; //konec sub-matice v bloku float Csub = 0; for (int a = aBegin, b = BLOCK_SIZE * bx; a <= aEnd; a += BLOCK_SIZE, b += BLOCK_SIZE * n) { __shared__ float As[BLOCK_SIZE][BLOCK_SIZE]; //sdílená pamˇ et’ __shared__ float Bs[BLOCK_SIZE][BLOCK_SIZE]; //v rámci bloku As[ty][tx] = A[a + n * ty + tx]; //každé vlákno naˇ cte jeden Bs[ty][tx] = B[b + n * ty + tx]; //element matice do pamˇ eti __syncthreads(); //synchronizace, sub-matice ve sdílené pamˇ eti for (int k = 0; k < BLOCK_SIZE; ++k) { //každé vlákno spoˇ cítá Csub += As[ty][k] * Bs[k][tx]; //dílˇ ci element sub-matice } __syncthreads(); } int c = n * BLOCK_SIZE * by + BLOCK_SIZE * bx; C[c + n * ty + tx] = Csub; //zápis výsledku do device pamˇ eti }



12 / 14

CUDA - pˇríklad - Násobení matic 7/8 ˇ CUDA zdrojových kódu. Zaˇclenení ˚

Pˇríklad - samostatný zdrojový soubor cuda_func.cu 1. Deklarace externí funkce. 1 2 3

extern "C" { //deklarace externí funkce void cuda_multiply(const int n, const float *A, const float *B, float *C); }

2. Vytvoˇrení zdrojového C++. 1

nvcc --cuda cuda_func.cu -o cuda_func.cu.cc

ˇ 3. Kompilace cuda_func.cu.cc bežným kompilátorem.



13 / 14

CUDA - pˇríklad - Násobení matic 8/8 ˇ doby výpoˇctu - cˇ as v ms Prub ˚ eh 18000

16000

naive naive transposition cuda

Computation time [ms]

14000

12000

10000

8000

6000

4000

2000

0 200

400

600

800

1000

1200

1400

Matrix size

N 112 208 304

Naive 2 11 35

•

Transp. 1 11 33

CUDA 81 82 84

N 704 1104 1264

Naive 1083 6360 9763

Transp. 405 1628 2485

CUDA 122 235 308

Matlab 7.6.0 (R2008a): n=1104; A=rand(n,n); B=rand(n,n); tic; C=A*B; toc Elapsed time is 0.224183 seconds.



14 / 14

GPGPU. Jan Faigl. Gerstnerova Laboratoř pro inteligentní rozhodování a řízení České vysoké učení technické v Praze

Recommend Documents