CUDA cudniejsze przykady Agenda CPU vs GPU Mnoenie

CUDA cudniejsze przyk|ady

Agenda: | CPU vs. GPU |Mnożenie macierzy – CPU | Mnożenie macierzy - GPU |Sploty

Macierze – CPU vs. GPU CPU: | Mnożenie wykonywane w kolejnych iteracjach pętli. | Przechodzimy przez pierwszy wiersz tabeli M i pierwszą kolumnę macierzy N, w pętli liczymy ich iloczyn skalarny i od razu zapisujemy go w macierzy wyjściowej. GPU: | paralelne wykonanie mnożenia macierzy właściwego | każdy wątek wpisuje jeden element do macierzy wynikowej Cdn. . .

CPU Istotny dla zrozumienia różnicy pomiędzy tradycyjnym mnożeniem macierzy, a analogicznymi obliczeniami na GPU jest fakt linearnego adresowania macierzy w pierwszym przypadku. Właśnie z linearnego adresowania macierzy wynika występujący w iteracji indeks postaci: – i*Width+k. (i – numer wiersza; Width – wymiar n macierzy; k-numer kolumny) Adresowanie linearne pokazano schematycznie na ilustracji:

CPU Jak widać, jest to podejście typowo sekwencyjne…

Coś ciekawszego, czyli: GPU

GPU - (C = AB) Wersja najprostsza: Jest to wersja niezoptymalizowana. Każdy half warp oblicza jeden rząd tile’a C, polegając przy tym na jednym rzędzie z A i całym tile’u z B. W każdej iteracji pętli wszystkie wątki w half warpie czytają tę samą wartość z pamięci globalnej. __global__ void simple. Multiply(float *a, float* b, float *c, int N) { int row = block. Idx. y * block. Dim. y + thread. Idx. y; int col = block. Idx. x * block. Dim. x + thread. Idx. x; float sum = 0. 0 f; for (int i = 0; i < TILE_DIM; i++) { sum += a[row*TILE_DIM+i] * b[i*N+col]; } c[row*N+col] = sum; }

Wersja 2: GPU - (C = AB) Pierwszym z możliwych ulepszeń jest wykorzystanie pamięci współdzielonej. W drugiej wersji algorytmu wczytujemy tile z A do pamięci współdzielonej. __global__ void coalesced. Multiply(float *a, float* b, float *c, int N) { __shared__ float a. Tile[TILE_DIM]; int row = block. Idx. y * block. Dim. y + thread. Idx. y; int col = block. Idx. x * block. Dim. x + thread. Idx. x; float sum = 0. 0 f; a. Tile[thread. Idx. y][thread. Idx. x] = a[row*TILE_DIM+thread. Idx. x]; for (int i = 0; i < TILE_DIM; i++) { sum += a. Tile[thread. Idx. y][i]* b[i*N+col]; } c[row*N+col] = sum; }

Wersja 3: GPU - (C = AB) Kolejnym możliwym ulepszeniem jest jednorazowe wczytywanie całego rzędu macierzy B do pamięci współdzielonej. __global__ void shared. ABMultiply(float *a, float* b, float *c, int N) { __shared__ float a. Tile[TILE_DIM], b. Tile[TILE_DIM]; int row = block. Idx. y * block. Dim. y + thread. Idx. y; int col = block. Idx. x * block. Dim. x + thread. Idx. x; float sum = 0. 0 f; a. Tile[thread. Idx. y][thread. Idx. x] = a[row*TILE_DIM+thread. Idx. x]; b. Tile[thread. Idx. y][thread. Idx. x] = b[thread. Idx. y*N+col]; __syncthreads(); for (int i = 0; i < TILE_DIM; i++) { sum += a. Tile[thread. Idx. y][i]* b. Tile[i][thread. Idx. x]; } c[row*N+col] = sum; }

Porównanie: Capability 1. 1 - NVIDIA Ge. Force GT 9600 M in a Mac. Book Pro Laptop, 4 multiprocessors, 32 cores Capability 1. 2 - NVIDIA Ge. Force GT 330 M in a Mac. Book Pro Laptop, 6 multiprocessors, 48 cores

Porównanie: Capability 1. 3 - NVIDIA Tesla C 1060 running in Earlham's cluster, 30 multiprocessors, 240 cores Capability 2. 0 - NVIDIA Tesla M 2070 at the Texas Advanced Computing Center, 14 multiprocessors, 448 cores Matrix Multiplication with CUDA | A basic introduction to the CUDA programming model Robert Hochberg

Porównanie: NVIDIA Ge. Force Optymalizacja GTX 280(1. 3) GTX 8800(1. 0) Bez optymalizacji 8. 7 GBps 0. 7 GBps __shared__ float a. Tile 14. 3 GBps 8. 2 GBps __shared__ float a. Tile, b. Tile 29. 7 GBps 15. 7 GBps Matrix Multiplication with CUDA | NVIDIA CUDA C Best Practices Guide

Sploty Splot znajduje szerokie zastosowanie w przetwarzaniu obrazów. Operacja splotu oblicza nową wartość piksela obrazu na podstawie wartości pikseli sąsiadujących. Przed zastosowaniem splotu: Po zastosowaniu splotu:

Sploty Simple box blur: maska: Przed zastosowaniem splotu: Po zastosowaniu splotu: (jak widać, maska ma efekt uśredniający)

Sploty Gaussian blur: Przed zastosowaniem splotu: maska: Po zastosowaniu splotu:

Sploty Naiwna implementacja: W najprostszej wersji implementacji splotu każdy blok wątków przetwarza jeden blok obrazu. Każdy wątek generuje na wyjściu jeden piksel.

Sploty Brutalna konfrontacja: (z rzeczywistością) Po zmodyfikowaniu naiwnego algorytmu zagadnienie zaczyna się komplikować… Uwzględnienie w algorytmie „otoczki”, niezbędnej do przeliczenia brzegowych pikseli powoduje, że wątki odpowiedzialne wcześniej za wczytanie „otaczających” pikseli będą bezczynne przez cały czas przeliczania maski.

Autorki | Urszula Jędrzejczak | Katarzyna Ostrowicz

Koniec Bibliografia: | „CUDA by Example: An Introduction to General-Purpose GPU Programming” Jason Sanders, Edward Kandrot | „Programming Massively Parallel Processors: A Hands-on Approach” David B. Kirk, Wen-mei W. Hwu | | Dokumentacja NVIDIA: CUDA C Best Practices Guide Version 3. 1 z 2010 -05 -28 | Image Convolution with CUDA Victor Podlozhnyuk[http: //developer. download. nvidia. com/compute/cuda/1. 1 - Beta/x 86_64_website/projects/convolution. Separable/doc/convolution. Separabl e. pdf] | Matrix Multiplication with CUDA | A basic introduction to the CUDA programming model. Robert Hochberg [http: //www. shodor. org/media/content//petascale/materials/UPModule s/matrix. Multiplication/module. Document. pdf] | http: //www. aishack. in/2010/08/image-convolutionexamples/