1/52
Vlastnosti CUDA Metriky algoritmů Měření
Metriky efektivity na GPU
Petr Holub
hopet@ics.muni.cz
Laboratoř pokročilých síťových technologií
PV197
2010–11–17
2/52
Vlastnosti CUDA Metriky algoritmů Měření
Přehled přednášky
Vlastnosti CUDA
Metriky algoritmů
Měření
3/52
Vlastnosti CUDA Metriky algoritmů Měření
Literatura
Park I. K., Singhal N., Lee M. H., Cho S., Kim C. W., “Design and
Performance of Evaluation of Image Processing Algorithms on GPUs,”
IEEE Transactions on Parallel and Distributed Systems, 2010
(zatím pouze v elektronické verzi)
Cope B., Cheung P. Y. K., Luk W., Howes L., “Performance Comparison
of Graphics Processors to Reconfigurable Logic: A Case Study”, IEEE
Transactions on Copmuters, vol. 59, no. 4, April 2010
Best Practices Guide – CUDA 2.2, 2009 http://developer.
download.nvidia.com/compute/cuda/2_3/toolkit/
docs/NVIDIA_CUDA_BestPracticesGuide_2.3.pdf
Wil Braithwaite, “The CUDA architecture: The Art of performance
optimization”, Siggraph 2009, http:
//developer.download.nvidia.com/presentations/
2009/SIGGRAPH/asia/6_cuda_optimization.pdf
4/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s pamětí
Práce s pamětí
◾ omezení pamětí
omezená šířka pásma mezi host a device
(cca 6 GB/s pro PCI-e x16 Gen2)
latence globální paměti (teoreticky
141 GB/s, cca 400–600 cyklů latence)
1107 × 106
takt paměti [Hz]
× (512/8)
Inteface paměti
× 2
DDR
/109 = 141,6 GB/s
(ev. 132 GB/s při dělení 10243)
omezená velikost sdílené paměti
omezený počet registrů
CPU
chipset
paměť
CPU
GPU
paměť
GPU
PCIe, 5 GB/s
PCIe,5–10GB/s
PCIe,60–80GB/s
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
5/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s pamětí
Práce s pamětí
◾ optimalizace
maximální využití sdílené paměti a
registrů
koalescentní přístup do globální paměti a
vyhnutí se partition campingu
překrývání výpočtů a přístupu k datům
využití asynchronních přenosů:
cudaMemcpyAsync(...);
přiměřené používání page-locked paměti:
cudaHostAlloc(...);
co největší přenosy host ↔ device
najednou
CPU
chipset
paměť
CPU
GPU
paměť
GPU
PCIe, 5 GB/s
PCIe,5–10GB/s
PCIe,60–80GB/s
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
6/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s pamětí
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
7/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Další triky s pamětí
◾ rozvrhnout, co cacheovat a co opakovaně počítat
◾ mapování OpenGL bufferu do adresního prostoru zařízení (device)
1. zaregistrujte si buffer pomocí CUDA-C
cudaGLRegisterBufferObject(GLuint buffObj);
2. namapujte zaregistrovaný buffer do globální paměti zařízení (vrátí adresu)
cudaGLMapBufferObject(void** devPtr,GLuint buffObj);
3. použijte adresu v kernelu
4. odmapujte buffer
cudaGLUnmapBufferObject(GLuint buffObj);
5. odregistrujte buffer
cudaGLUnregisterBufferObject(GLuint buffObj); (potřeba pouze
pokud je buffer cíl rendrování)
6. použijte buffer v OpenGL
může pomoci odstranit přenosy host ↔ device
automatické DMA mezi kartami Tesla a Quadro (momentálně přes host)
vykreslování z pixel buffer object pomocí glDrawPixels nebo
glTexImage2D
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
8/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Textura generovaná pomocí CUDA:
// setup code:
2 cudaGLRegisterBufferObject(pbo);
// CUDA texture generation code:
4 unsigned char *d_buffer;
cudaGLMapBufferObject((void**)&d_buffer, pbo);
6 prep_texture_kernel<<<...>>>(d_buffer);
cudaGLUnmapBufferObject(pbo);
8 // OpenGL rendering code:
glBindBuffer(GL_PIXEL_UNPACK_BUFFER_ARB, pbo);
10 glBindTexture(GL_TEXTURE_2D, tex);
glTexSubImage2D(GL_TEXTURE_2D,0,0,0,256,256,GL_BGRA,GL_UNSIGNED_BYTE,0);
Zpracování snímku pomocí CUDA:
// OpenGL rendering code:
2 // ...
// CUDA post-processing code:
4 unsigned char *d_buffer;
cudaGLRegisterBufferObject(pbo);
6 cudaGLMapBufferObject((void**)&d_buffer, pbo);
post_process_kernel<<<...>>>(d_buffer);
8 cudaGLUnmapBufferObject(pbo);
cudaGLUnRegisterBufferObject(pbo);
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
9/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Další triky s pamětí
◾ Write-Combining
cudaHostAlloc((void**)&h_data, num_bytes,
cudaHostAllocWriteCombined);
paměť není cacheovaná ani cache koherentní
PCI nedělá snooping
podle CUDA 2.2 Pinned Memory APIs
(http://www.fcsc.es/download/Archivo%20Cursos/CUDA_
Unileon_2009/CUDA2.2PinnedMemoryAPIs.pdf) může
poskytnout až o 40 % větší výkon
může zvýšit i výkon pro zápis procesorem (Write Combining Memory
Implementation Guidelines, http://download.intel.com/
design/PentiumII/applnots/24442201.pdf): agregací zápisu,
obcházením L1/L2 cache
problematické čtení – potřeba používat paměťové bariéry před čtením
hodnot (_mm_sfence na Linuxu, _WriteBarrier na Windows, u SSE4 lze
pro čtení použít instrukci MOVNTDQA); bariéry provádí CUDA driver a jsou
relativně pomalé
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
10/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Události
◾ synchronizace omezená na specifickou událost
cudaEvent_t HtoDdone;
2 cudaEventCreate(&HtoDdone,0);
cudaMemcpyAsync(d_dest,h_source,bytes,cudaMemcpyHostToDevice,0);
4 cudaEventRecord(HtoDdone);
myKernel<<<grid,block>>>(...);
6 cudaMemcpyAsync(d_dest,h_source,bytes,cudaMemcpyDeviceToHost,0);
// cpu can do stuff here
8 cudaEventSynchronize(HtoDdone);
// waits just for everything before
10 // cudaEventRecord(HtoDdone)
// to complete and then runs
12 //
// The first memory copy is done,
14 // so the memory at source could be
// used again by the CPU
16 cudaThreadSynchronize();
// waits for everyting on GPU to finish, then returns
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
11/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Streamy
◾ identifikace při asynchronních přenosech
◾ cudaStream_t, cudaStreamCreate()
◾ fronty se plní v pořadí vykonávání kódu
CudaMemcpyAsync(A1...,StreamA);
2 KernelA1<<<...,StreamA>>>();
KernelA2<<<...,StreamA>>>();
4 KernelA3<<<...,StreamA>>>();
CudaMemcpyAsync(A2...,StreamA);
6 CudaMemcpyAsync(B1...,StreamB);
CudaMemcpyAsync(B2...,StreamB);
8 KernelB1<<<...,StreamB>>>();
CudaMemcpyAsync(B3...,StreamB);
10 CudaMemcpyAsync(B4...,StreamB);
Špatně – blokováno vykonávání kernelu B1 kvůli kopírování výsledků
A2
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
12/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Streamy
◾ identifikace při asynchronních přenosech
◾ cudaStream_t, cudaStreamCreate()
◾ fronty se plní v pořadí vykonávání kódu
CudaMemcpyAsync(A1...,StreamA);
2 KernelA1<<<...,StreamA>>>();
KernelA2<<<...,StreamA>>>();
4 KernelA3<<<...,StreamA>>>();
CudaMemcpyAsync(B1...,StreamB);
6 CudaMemcpyAsync(B2...,StreamB);
KernelB1<<<...,StreamB>>>();
8 CudaMemcpyAsync(A2...,StreamA);
CudaMemcpyAsync(B2...,StreamB);
10 CudaMemcpyAsync(B2...,StreamB);
OK
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
13/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Streamy
◾ vícecestné překryvy (3-way, 4-way)
Wil Braithwaite, “The CUDA architecture: The Art
of performance optimization”
◾ Fermi obsahuje druhou kopírovací jednotku – zajímavé pro 4-cestné
přenosy
Způsoby čekání na synchronní události
◾ cudaError_t cudaSetDeviceFlags (int flags)
◾ parametry:
cudaDeviceScheduleSpin – minimalizuje latenci, CPU vlákno aktivně spinuje
cudaDeviceScheduleYield – umožňuje co nejefektivnější běh vlákna na CPU
cudaDeviceBlockingSync – CPU vlákno se při čekání zablokuje pomocí
synchronizačního primitiva
cudaDeviceScheduleAuto – default, heuristika: buď aktivně spinuje nebo provede
yield, podle počtu aktivních CUDA kontextů a počtu logických procesorů v systému
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
14/52
Vlastnosti CUDA Metriky algoritmů Měření
Pokročilá práce s pamětí
Zero-copy mapování paměti
◾ přímý přístup k datům v paměti CPU
◾ schopnost tohoto se testuje pomocí pole canMapHostMemory dotazu
cudaDeviceProp
◾ automatický přenos dat po PCIe dle potřeby
◾ relativně pomalé
pro jednorázově použitá malá data
pokud dokáže vysoký podíl výpočtu maskovat latenci
◾ zajímavé zejména v kombinaci s integrovanými kartami – integrated
pole cudaDeviceProp
na UMA architektuře odpadá přenos úplně
Nvidia ION
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
15/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Vlákna a multiprocesory
◾ potřebujeme využít výpočetní výkon karty
◾ kolik warpů potřebujeme k maskování latence globální paměti?
řekněme, že potřebujeme 100 aritmetických instrukcí k maskování latence
(400 taktů latence / 4 takty na instrukci)
řekněme, že máme 8 aritmetických instrukcí (8 × 4 takty) na 1 přístup do
globální paměti (400 taktů latence)
100/8 ≈ 13 warpů
◾ kolik warpů potřebujeme k maskování read-after-write latence registrů?
latence je cca 24 cyklů
24/4 ≈ 6 warpů
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
16/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Vlákna a multiprocesory
◾ obsazení (occupancy)
O =
# warpů běžících na MP v daný okamžik
maximální # souběžných warpů
c. c. 1.2: maximálně 32 warpů
c. c. 1.1: maximálně 24 warpů
13/32 = 40 % obsazení pro maskování latence globální paměti (c. c. 1.2)
6/32 = 18,75 % obsazení pro maskování latence registrů (c. c. 1.2)
◾ více vláken
?
= větší výkon
záleží na zdrojích požadovaných vláknem
záleží na uspořádání vláken
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
17/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Limity multiprocesoru:
◾ počty registrů: 16384 (64 kB) nebo jen 8192 (32 kB) na starším železe
◾ sdílená paměť: 16 kB
◾ maximální počet warpů: 32
◾ maximální počet bloků: 8
Pozor na dostatečný počet vláken/warpů v bloku (≥ 96)
Pozor na limity v počtu potřebných registrů
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
18/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Příklad limitů pro 8132 registrů a 24 warpů:
◾ 10 registrů na vlákno, 256 vláken na blok
každý blok použije 2650 registrů ⇒ mohou běžet 3 bloky (7680
registrů)
256 × 3/32 = 24 warpů může běžet současně ⇒ může dosáhnout
100 % využití
◾ 17 registrů na vlákno, 256 vláken na blok
každý blok použije 4352 registrů ⇒ může běžet 1 blok (4352 registrů)
256 × 1/32 = 8 warpů může běžet současně ⇒ může dosáhnout jen
33 % využití
◾ 17 registrů na vlákno, 128 vláken na blok
každý blok použije 2176 registrů ⇒ mohou běžet 3 bloky (6528
registrů)
128 × 3/32 = 12 warpů může běžet současně ⇒ může dosáhnout 50 %
využití
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
19/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Určení používaných zdrojů
◾ přeložíme s -cubin
◾ výsledný .cubin soubor obsahuje
architecture {sm_10}
2 abiversion {0}
modname {cubin}
4 code {
name = MyKernel
6 lmem = 0 // lokalni pamet per blok
smem = 68 // sdilena pamet per blok
8 reg = 20 // pocet registru per vlakno
bar = 0
10 bincode {
0xa0004205 0x04200780 0x40024c09 0x00200780
◾ nebo použijeme –ptxas-options=-v
ptxas info : Used 4 registers, 60+56 bytes lmem, 44+40 bytes
smem, 20 bytes cmem[1], 12 bytes cmem[14]
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
20/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Omezení tlaku na registry
◾ kompilátor se snaží počet registrů minimalizovat
◾ -maxrregcount=<N> umožňuje nastavit požadovaný maximální počet
registrů na kernel
◾ přetečení do lokální paměti způsobí zpomalení
Heuristiky pro velikost mřížky
◾ # bloků > # multiprocesorů
aby všechny multiprocesory měly alespoň jeden blok k vykonávání
◾ # bloků / # multiprocesorů > 2
na jednom procesoru může běžet více bloků
bloky nečekající v __syncthreads(); udržují zátěž hardware
◾ # bloků > 1000
rezerva do budoucnosti přes několik generací karet
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
21/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
Heuristiky na velikost bloku
◾ čím více vláken v bloku, tím méně může jedno vlákno používat registrů
◾ čím více vláken v bloku, tím hlubší pipeline a lepší maskování latence
◾ počty vláken v bloku by měly být násobky 64
minimalizace bank konfliktů na registrech
rozmné: 192 nebo 256
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
22/52
Vlastnosti CUDA Metriky algoritmů Měření
Shrnutí práce s procesory
CUDA Occupancy Calculator
http://developer.download.nvidia.com/compute/cuda/CUDA_Occupancy_
calculator.xls
Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization”
23/52
Vlastnosti CUDA Metriky algoritmů Měření
Metriky z pohledu obrazových algoritmů
1. Paralelní podíl
2. Poměr mezi operacemi v plovoucí čárce a přístupů do globální
paměti
3. Počet operací v plovoucí čárce na pixel
4. Počet přístupů do paměti na pixel
5. Míra větvení
6. Závislost úloh
Srovnávané algoritmy implementovány na CPU
◾ výpočet metrik před implementací na CUDA
Zdroj: Park et al., “Design and Performance Evaluation of Image Processing Algorithms on GPUs”
24/52
Vlastnosti CUDA Metriky algoritmů Měření
Metriky z pohledu obrazových algoritmů
Paralelní podíl
◾ odpovídá Amdahlovu zákonu se všemi důsledky
◾ při složeném algoritmu odpovídá poměrům částí
Poměr mezi operacemi v plovoucí čárce a přístupů do globální
paměti
◾ jakékoli zápisy (sdílená, lokální, globální paměť)
◾ skrývání latence překrýváním výpočtem
Počet operací v plovoucí čárce (FP) na pixel
◾ při výpočtech v plovoucí čárce překonávají GPU cca 20× CPU
◾ obrazové zpracování zahrnuje typicky hodně operací v plovoucí čárce
◾ nepřímo koreluje i s přístupy do paměti
Počet přístupů do paměti na pixel
◾ GPU mají přibližně 10× větší kapacitu přístupu do paměti než CPU
◾ per pixel charakteristika umožňuje často využití sdílené paměti
◾ problémy se sekvenčními přístupy do paměti u některých algoritmů –
omezují paralelismus
Zdroj: Park et al., “Design and Performance Evaluation of Image Processing Algorithms on GPUs”
25/52
Vlastnosti CUDA Metriky algoritmů Měření
Metriky z pohledu obrazových algoritmů
Míra větvení
◾ větvení přes if, switch, do, for, while
◾ algoritmy pro zpracování obrázků často používají větvení na základě
výsledku bitových operací
◾ počítá se na základě rozptylu délky běhu jednotlivých vláken
◾ berou se vlákna po blocích velikosti 32
Závislost úloh
◾ při zpracování obrázků se závislosti většinou řeší sekvenčním
spouštěním CUDA kernelů
◾ sleduje se počet bariér
Zdroj: Park et al., “Design and Performance Evaluation of Image Processing Algorithms on GPUs”
26/52
Vlastnosti CUDA Metriky algoritmů Měření
Relativní důležitost metrik
Paralelní podíl > Větvení > FP operace per pixel >
přístupy do paměti per pixel >
poměr mezi FP operacemi a přístupy do paměti > závislost úloh
1. Amdahl limituje vše
2. omezení SIMT modelu
3. algoritmy na zpracování obrazu mají obecně hodně FP operací
4. poměr je závislý na předchozích dvou hodnotách (proč ho zavádět?)
5. závislost úloh určuje obtížnost implementace (opravdu?)
27/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Stereo srovnávání obrázků (Multiview Stereo Matching, MVS)
28/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Stereo srovnávání obrázků (Multiview Stereo Matching, MVS)
◾ popis algoritmu
vstup:
zkalibrované obrazy I = I0, . . . ,IN−1
projekční matice P = P0, . . . ,PN−1
výstup:
3D body X = X0, . . . ,XM−1
porovnávání lokálních oken mezi jednotlivými obrazy
hloubka se pro bod (x,y) v referenčním obrazu určí přeložením příslušné
oblasti z Ii na oblasti v obraz Ii−1
bod s minimálním součtem absolutních odchylek (SAD) a normalizovanou
cross-korelací (NCC) je uložen jako best-match
opakujeme pro ostatní sousedící obrazy (např. Ii+1)
bod je korektně určen, pokud počet best-match je nad hranici
MIN_COUNT
opakujeme pro všechny body a všechny referenční obrazy
◾ složitost: O(N2
WHL)
N ...počet vstupních obrazů, W resp. H ...vodorovné resp. svislé
rozlišení, L ...velikost ohraničujícího boxu
29/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Stereo srovnávání obrázků (Multiview Stereo Matching, MVS)
◾ mapování na GPU
porovnávání lokálních oken je nezávislé – dobře mapuje na GPU
vlákno ∼ pixel, tj. W × H vláken
O(N) volání kernelu pro výpočet hloubky pro jeden referenční obraz
O(L) volání uvnitř kernelu
obrázky se nakopírují do globální paměti
koeficienty lokálního okna jdou do sdílené paměti – častý přístup
složitost: O( N2
WHL
Tmax
)
Tmax...maximální počet vláken na GPU (např. 12288 na G80)
30/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Získávání lineárních charakteristik (Linear Feature Extraction)
◾ rozpoznávací aplikace: budovy, silniční pruhy, ...
Y.
T. Zhou, “Linear Feature Extraction Based on an AR Model Edge Detector”
http://ieeexplore.ieee.org/stamp/stamp.jsp?arnumber=01169687
31/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Získávání lineárních charakteristik (Linear Feature Extraction)
◾ popis algoritmu
vstup: obrázek
výstup: obrázek s detekovanými hranami
použijeme algoritmus Nevatia-Babu s Cannyho metodou detekce hran
1. detekce hran
2. zúžení hran
3. vytvoření řetízků hran na základě spojitosti v 8 směrech
4. fitting čar na řetízky
5. pokud segment čáry dává větší chybu než hraniční hodnota, je čára
rozdělena na 2
6. poslední krok iterativně opakujeme
32/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Získávání lineárních charakteristik (Linear Feature Extraction)
◾ mapování na GPU
6 per-pixel kernelů
1. Cannyho detekce hran jako per-pixel filtr
2. klasifikace pixelů podle okolí 3 × 3 – Initialization
3. nalezení co nejdelších souvislých částí, určení počátečního a koncového
bodu – Linking
4. výpočet odchylky od úsečky proložené mezi počáteční a koncový bod –
Fitting
5. pokud odchylka přesahuje Dmax, rozdělí se souvislý řetízek na dva
6. iterujeme poslední 3 kroky, dokud vznikají nové segmenty
33/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Získávání lineárních charakteristik (Linear Feature Extraction)
◾ mapování na GPU
počítá se pro všechny pixely – i ty, co hrany neobsahují ⇒ omezení
redundantní hledání nejdelších souvislých čar
významný podíl over-computation
34/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
JPEG2000 – diskrétní vlnková transformace (DWT) a EBCOT
DWT /
Quantization
Output Bit Stream
Formating
Context
Modeling
CX
D
Arithmetic
Encoding
Tier-1
Tier-2
Arithmetic
Encoding
(10 %)
Context
Modeling
(51 %)
DWT (32 %)
Rest (7 %)
Matela, Rusˇnák, Holub, “GPU-Based Sample-Parallel Context Modelling for EBCOT in
JPEG2000”, MEMICS2010
DWT: per pixel, coefficient lifting
EBCOT Tier-1: problém se sériovou definicí algoritmu
◾ uvnitř bloku (code block) jsou složité závislosti
◾ algoritmus je definován tak, blokem prochází postupně a aktualizuje
stav
◾ mezi bloky závislosti nejsou
35/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
JPEG2000 – diskrétní vlnková transformace (DWT) a EBCOT
◾ mapování na GPU
DWT: per pixel
EBCOT Tier-1: per blok
36/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Nefotorealistický rendering
◾ Cartoon-style NPR
◾ Oily-style NPR
37/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Nefotorealistický rendering
◾ popis algoritmu: Cartoon-style NPR
bilaterální filtrování (plochy)
Cannyho detekce hran (čáry)
přeložení čar přes plochy
◾ mapování na GPU: Cartoon-style NPR
per pixel pro všechny operace
vstupní obrázky jsou uloženy ve 2D texturové paměti
38/52
Vlastnosti CUDA Metriky algoritmů Měření
Studované pilotní algoritmy
Nefotorealistický rendering
◾ popis algoritmu: Oily-style NPR
rozdíl původního obrazu a gaussovsky rozostřeného obrazu per pixel (míra
rozostření udává tloušťku štětce)
tah se generuje, pokud oblast v součtu rozdílů dosahuje nad stanovenou
hranici
tah začíná od lokálního maxima a vede po gradientu
tah se uloží pouze je-li dost dlouhý
uložené tahy se ,,obtáhnou štětcem‘‘ dané tloušťky
proces začíná s čistým pozadím a opakuje se od nejtlustšího štětce k
nejtenčímu
◾ mapování na GPU: Oily-style NPR
paralelizace per tah: problém s překryvem tahů
paralelizace per pixel: problém jak vybrat poředí tahů
reformuluace/heuristika: světlejší štětec se použije později než tmavší
⇒ pro daný pixel vybereme nejsvětlejší barvu překrývajících se tahů
musíme hledat, kterými tahy bude pixel ovlivněn (maximální vzdálenost
od trajektorie tahu)
vstupní obrázky jsou uloženy ve 2D texturové paměti
39/52
Vlastnosti CUDA Metriky algoritmů Měření
Přehled pilotních algoritmů
40/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
charakteristika algoritmů
◾ MVS kompenzuje malý poměr FLOP/mem vysokou mírou paralelismu
◾ DWT kompenzuje málo FLOP/pixel vysokou mírou paralelismu
◾ LFE má problém s nízkou mírou paralelismu
41/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
charakteristika algoritmů
◾ task depenednecy: implementace bude stát hodně úsilí (schoval se
nám díky zvolenému přístupu EBCOT Tier-1)
⇒ Oily-Style NPR má mnoho iterací, během nichž aktualizuje buffer
◾ EBCOT Tier-1 má problém s větvením, nižším paralelismem a malým
počtem FLOP/pixel
42/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
benchmarky CUDA
◾ DWT: agregace dlaždic sníží potřebný počet registrů ( ⇒ GPU
occupancy)
43/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
benchmarky CUDA
◾ Data Dependency: celkový počet volání __syncthreads(); v rámci
bloku vláken
44/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
zrychlení
45/52
Vlastnosti CUDA Metriky algoritmů Měření
Výsledky pilotních algoritmů
zrychlení
46/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
velikost obrázku čas běhu
640 × 480 124,12983930928
1280 × 720 539,98450298239
1920 × 1080 1529,02398429008
4096 × 2160 10210,09238488922
47/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
velikost obrázku čas běhu
640 × 480 124,12983930928
1280 × 720 539,98450298239
1920 × 1080 1529,02398429008
4096 × 2160 10210,09238488922
TAKHLE NE!
Co je špatně? ...skoro všechno
◾ jednotky
◾ jak byl sestaven systém pro měření
◾ jak bylo měření provedeno
◾ jak bylo měření vyhodnoceno
◾ opravdu jsme měřili až na 16 platných míst?
48/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
Návrh experimentu
◾ co chceme změřit? – formulujeme hypotézu, kterou ověřujeme
◾ máme měření nastavené tak, abychom změřili skutečně to, co chceme?
◾ neměříme např. místo výpočetního výkonu propustnost nějaké
sběrnice?
Protokolování měření
◾ podklady pro zajištění reprodukovatelnosti měření
◾ příklad s měřením AGC po půl roce...
49/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
Statistické zpracování expirimentu (ve velmi malé, nedostatečné, a občas až za hranice přijatelnosti
zjednodušené kostce)
◾ chyby v měření: systematické, náhodné
◾ opakování měření: eliminace náhodných chyb
50/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
Statistické zpracování expirimentu (ve velmi malé, nedostatečné, a občas až za hranice přijatelnosti
zjednodušené kostce)
◾ náhodná chyba měření (směrodatná odchylka průměru) klesá s
√
n, kde
n je počet měření
běžný počet opakování je 5–15
◾ použít průměr, medián nebo minimální hodnotu?
jaká je distribuce výsledků?
co chceme výsledkem říci?
◾ stačí nám jedno číslo jako výsledek?
51/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
Statistické zpracování expirimentu (ve velmi malé, nedostatečné, a občas až za hranice přijatelnosti
zjednodušené kostce)
◾ při intervalu důvěryhodnosti 95 % a normálním rozdělení se udává
x ± 1,96σx [jednotka]
kde průměr x
x =
N
∑
i=1
xi
N
směrodatná neboli standardní odchylka σx
σx =
1
N − 1
N
∑
i=1
(xi − x)2 =
1
N − 1
(
N
∑
i=1
x2
i − Nx2
)
směrodatná odchylka aritmetického průměru σx
σx =
σx
√
N
52/52
Vlastnosti CUDA Metriky algoritmů Měření
Pár poznámek k měření...
Statistické zpracování expirimentu (ve velmi malé, nedostatečné, a občas až za hranice přijatelnosti
zjednodušené kostce)
◾ chybu zaokrouhlujeme na 1, výjimečně 2 desetinná místa a
zaokrouhlení průměru tomu přizpůsobíme
◾ můžeme mít ne-normální rozdělení (např. Poissonovo)
◾ propagace chyb
◾ můžeme potřebovat složitější statistické zpracování
◾ omlouvám se statistice za příliš zjednodušený a místy ne moc korektní
výklad :-)
Pořádné materiály o experimentech vydají na přednášku a cvika