1 / 19

CUDA

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.

irene-dalton
Download Presentation

CUDA

An Image/Link below is provided (as is) to download presentation Download Policy: Content on the Website is provided to you AS IS for your information and personal use and may not be sold / licensed / shared on other websites without getting consent from its author. Content is provided to you AS IS for your information and personal use only. Download presentation by click this link. While downloading, if for some reason you are not able to download a presentation, the publisher may have deleted the file from their server. During download, if you can't get a presentation, the file might be deleted by the publisher.

E N D

Presentation Transcript

  1. CUDA cudniejsze przyk|ady

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

  3. 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 właściwego mnożenia macierzy • każdy wątek wpisuje jeden element do macierzy wynikowej Cdn...

  4. 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:

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

  6. Coś ciekawszego, czyli: GPU

  7. GPU - (C = AB) • __global__ voidsimpleMultiply(float *a, float* b, float *c, intN) • { • introw = blockIdx.y * blockDim.y + threadIdx.y; • intcol = blockIdx.x * blockDim.x + threadIdx.x; • float sum = 0.0f; • for (int i = 0; i < TILE_DIM; i++) { • sum += a[row*TILE_DIM+i] * b[i*N+col]; • } • c[row*N+col] = sum; • } 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.

  8. GPU - (C = AB) • __global__ voidcoalescedMultiply(float *a, float* b, float *c, • int N) • { • __shared__ floataTile[TILE_DIM][TILE_DIM]; • introw = blockIdx.y * blockDim.y + threadIdx.y; • int col = blockIdx.x * blockDim.x + threadIdx.x; • float sum = 0.0f; • aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x]; • for (int i = 0; i < TILE_DIM; i++) { • sum += aTile[threadIdx.y][i]* b[i*N+col]; • } • c[row*N+col] = sum; • } Wersja 2: Pierwszym z możliwych ulepszeń jest wykorzystanie pamięci współdzielonej. W drugiej wersji algorytmu wczytujemy tile z A do pamięci współdzielonej.

  9. GPU - (C = AB) __global__ voidsharedABMultiply(float *a, float* b, float *c, int N) { __shared__ floataTile[TILE_DIM][TILE_DIM], bTile[TILE_DIM][TILE_DIM]; introw = blockIdx.y * blockDim.y + threadIdx.y; int col = blockIdx.x * blockDim.x + threadIdx.x; floatsum = 0.0f; aTile[threadIdx.y][threadIdx.x] = a[row*TILE_DIM+threadIdx.x]; bTile[threadIdx.y][threadIdx.x] = b[threadIdx.y*N+col]; __syncthreads(); for (int i = 0; i < TILE_DIM; i++) { sum += aTile[threadIdx.y][i]* bTile[i][threadIdx.x]; } c[row*N+col] = sum; } Wersja 3: Kolejnym możliwym ulepszeniem jest jednorazowe wczytywanie całego rzędu macierzy B do pamięci współdzielonej.

  10. Porównanie: Capability 1.1 - NVIDIA GeForce GT 9600M in a MacBook Pro Laptop, 4 multiprocessors, 32 cores Capability 1.2 - NVIDIA GeForce GT 330M in a MacBook Pro Laptop, 6 multiprocessors, 48 cores

  11. Porównanie: Capability 1.3 - NVIDIA Tesla C1060 running in Earlham's cluster, 30 multiprocessors, 240 cores Capability 2.0 - NVIDIA Tesla M2070 at the Texas Advanced Computing Center, 14 multiprocessors, 448 cores Matrix Multiplication with CUDA | A basic introduction to the CUDA programming model Robert Hochberg

  12. Porównanie: Matrix Multiplication with CUDA | NVIDIA CUDA C Best Practices Guide

  13. 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:

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

  15. Sploty Gaussianblur: maska: Przed zastosowaniem splotu: Po zastosowaniu splotu:

  16. 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.

  17. 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.

  18. Autorki • Urszula Jędrzejczak • Katarzyna Ostrowicz

  19. 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/x86_64_website/projects/convolutionSeparable/doc/convolutionSeparable.pdf] • Matrix Multiplication with CUDA | A basic introduction tothe CUDA programming modelRobert Hochberg[http://www.shodor.org/media/content//petascale/materials/UPModules/matrixMultiplication/moduleDocument.pdf] • http://www.aishack.in/2010/08/image-convolution-examples/

More Related