Orvosi kpfeldolgozsi algoritmusok prhuzamos megvalstsa CUDA technolgival Nikhzy

Orvosi képfeldolgozási algoritmusok párhuzamos megvalósítása CUDA technológiával Nikházy László Ureczky Bálint Konzulens: dr. Horváth Gábor

I. GPGPU és a CUDA 2

GPGPU General Purpose Computation on GPUs 3

CPU vs GPU 4

NVIDIA® CUDA™ Compute Unified Device Architecture GPGPU elősegítése • általános programozhatóság • shader-nyelvek • számítógépes grafika • skálázhatóság • szál-hierarchia • memória-hierarchia 5

A CUDA Architektúra 6

CUDA támogatottság • C/C++ § C Runtime for CUDA • . NET nyelvek § CUDA. NET – http: //www. gass- ltd. co. il/en/products/cuda. net • Java § Ja. Cuda – http: //jacuda. wiki. sourceforge. net • Python, Fortran, … § Windows / Linux / Mac. OS 7

Kernelfüggvények 8

Szálhierarchia 9

Memóriahierarchia 10

NVIDIA Ge. Force 9800 GTX+ • 16 multiprocesszor • 8 szálprocesszor / MP • 128 szálprocesszor • Memória: • 512 MB GDDR 3 • 256 -bit busz • 1100 MHz • 70. 4 GB/s • 1. 1 -es számítási kapacitás 11

1. 0 - ás számítási kapacitás: • rácsdimenziók : 65535 , 1 • blokkdimenziók : 512 , 64 • szál/proc : 512 • konstans memória : 64 KB • MP : 8 KB • regiszter/MP : 8192 • aktív blokk : 8 • atív szál : 768 1. 1 - es számítási kapacitás: • 32 -bites atomi függvények 12

II. CUDA programozás 13

Program felépítés • A hoszton és az eszközön futó kód együtt • C nyelvi kiterjesztések • Szétválasztás fordítási időben • Ajánlott fejlesztőkörnyezet: Microsoft Visual Studio 14

Nyelvi kiterjesztések - függvények • __host__ – a hoszton fut, csak a hosztról lehet hívni • __device__ – az eszközön fut, csak az eszközről hívható • __global__ – – kernel függvény a párhuzamosítás eszköze az eszközön fut, csak a hosztról lehet hívni Pl. : deklaráció: __global__ void func(float* parameter); meghívás: func<<<dim. Grid, dim. Block>>>(parameter); 15

Nyelvi kiterjesztések - változók • __device__ – az eszköz globális memóriájában foglal helyet • __shared__ – egy szál blokk közös memóriájában van – csak a blokkban lévő szálakból érhető el • __constant__ – az eszköz konstans memóriaterületén helyezkedik el – minden szálból, ill. a hosztról is elérhető • Beépített változók (kernel függvényekben) – grid. Dim, block. Dim – block. Idx, thread. Idx 16

Példa – mátrix összeadás __global__ void mat. Add(float A[N][N], float B[N][N], float C[N][N]) { int i = block. Idx. x * block. Dim. x + thread. Idx. x; int j = block. Idx. y * block. Dim. y + thread. Idx. y; if (i < N && j < N) C[i][j] = A[i][j] + B[i][j]; } int main() { // kernel hivas dim 3 dim. Block(16, 16); dim 3 dim. Grid((N + dim. Block. x – 1) / dim. Block. x, (N + dim. Block. y – 1) / dim. Block. y); mat. Add<<<dim. Grid, dim. Block>>>(A, B, C); } 17

Memória menedzsment + Textúra memória: cache-elt • foglalás, felszabadítás: – cuda. Malloc(void** ptr, size_t nbytes); – cuda. Free(void* ptr); • másolás: – cuda. Memcpy(void* dst, void* src, size_t nbytes, enum cuda. Memcpy. Kind direction); 18

Szinkronizáció • Hoszt - eszköz – kernel hívás aszinkron – memória másolás szinkron • Szálak között – ha több szál írja/olvassa ugyanazt a memóriaterületet, a sorrend nem determinisztikus – megoldás: __syncthreads(); • szinkronizációs pont: csak akkor halad tovább a végrehajtás, ha a blokkban minden szál eljutott eddig az utasításig 19

CUDA program általános menete 1. Adatok beolvasása a hoszton 2. Adatok másolása a hosztról az eszköz globális memóriájába 3. Kernel függvény hívása I. adatok másolása a globális memóriából a közös memóriába II. számítások III. eredmények visszaírása a közös memóriából a globális memóriába párhuzamos 4. Eredmény másolása az eszközről a hoszt memóriájába 5. Eredmény kiírása 20

III Párhuzamosan megvalósított algoritmusok

Kulcscsontkeresés 22

Kulcscsontkeresés § Előfeldolgozás (gyors) § Radon transzformáció (~1 -2 mp) kulcscsont főtengelyének meghatározása § Snake algoritmus (~2 -3 mp): Kulcscsont körvonalának iteratív meghatározása § Kulcscsont-eltűntetés… 23

Előfeldolgozás 1. Bejövő kép 2. Gauss szűrés 3. Gradiens számítás 4. Bejövő képpel pixelenkénti szorzás 5. Kulcscsont iránya szerinti szurés 6. Maximális értékek „elfojtása” Kimeneti kép 1 3 2 6 4 5 5 24

Radon-transzformáció (2 D-s, vonal szerinti Randon transzformáció) 1 0 -1 -90° -45° 0° 45° 90° 25

Mátrix konvolúció 26

Párhuzamos megvalósítás • Az eredmény mátrix minden egyes eleméhez külön szál • Egy szálon belül az összeg kiszámítása for ciklussal • Blokk méret: 16 x 16 • Textúra memória használata 27

Tesztelés: Prewitt operátor 28

FFT • DFT: • Radix-2 FFT: A A+B B A-B 29

Butterfly diagram x(0) x(4) X(0) x(2) x(6) X(1) W 80 X(2) W 82 X(3) x(1) W 80 X(4) x(5) W 81 X(5) W 82 X(6) W 83 X(7) W 80 x(3) x(7) W 80 W 82 30

Párhuzamos megvalósítás • Annyi szál, ahány pillangó egy szinten (N/2) • Minden iterációban minden szál egy pillangó műveletet végez • Közös memória használata • __syncthreads() iterációnként • 2 D FFT: 1 D FFT-k sorozata horizontálisan, majd vertikálisan 31

Teszt 32

Kitekintés • 100 -szoros gyorsítás az elvi határ • Számításintenzív feladatoknál (ahol sok egyszerű számolási műveletet végzünk egy adattal) ez jó közelítéssel elérhető • Memóriaintenzív feladatoknál 20 -30 -szoros gyorsítás reális • A tüdőkörvonal meghatározás inkább az utóbbi jellegű 33