1/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Efektivita na GPU Petr Holub hopet@ics.muni.cz SITOLA PV197 2011–11–28 2/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Přehled přednášky Vlastnosti CUDA Metriky algoritmů JPEG2000 3/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Shrnutí práce s pamětí Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization” 7/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 11/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 12/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 13/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Shrnutí práce s procesory Limity multiprocesoru: ◾ počty registrů: 16384 (64 kB) nebo jen 8192 (32 kB) na starším hardware ◾ 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” 14/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 15/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 16/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Shrnutí práce s procesory Omezení tlaku na registry ◾ kompilátor se snaží počet registrů minimalizovat ◾ -maxrregcount= 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” 17/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 rozumné: 192 nebo 256 Zdroj: Wil Braithwaite, “The CUDA architecture: The Art of performance optimization” 18/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 19/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 20/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 paměti ◾ jakékoli přístupy (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” 21/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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” 22/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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?) 23/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Studované pilotní algoritmy Stereo srovnávání obrázků (Multiview Stereo Matching, MVS) 24/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 25/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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) 26/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 27/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 28/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 29/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 30/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Studované pilotní algoritmy Nefotorealistický rendering ◾ Cartoon-style NPR ◾ Oily-style NPR 31/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 32/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 33/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Přehled pilotních algoritmů 34/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 35/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 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 36/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Výsledky pilotních algoritmů benchmarky CUDA ◾ DWT: agregace dlaždic sníží potřebný počet registrů ( ⇒ GPU occupancy) 37/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Výsledky pilotních algoritmů benchmarky CUDA ◾ Data Dependency: celkový počet volání __syncthreads(); v rámci bloku vláken 38/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Výsledky pilotních algoritmů zrychlení 39/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Výsledky pilotních algoritmů zrychlení 40/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 JPEG2000 – přehled procesu Source image Color transforms DWT and quanti- zation Context modeling Arithmetic coding EBCOT Tier-1 Output stream formating EBCOT Tier-2 Output JPEG2000 Literatura z naší zahrádky: :-)) ◾ MATELA, Jiří. GPU-Based DWT Acceleration for JPEG2000. In Annual Doctoral Workshop on Mathematical and Engineering Methods in Computer Science. Brno : NOVPRESS s.r.o., 2009. od s. 136-143, 8 s. ISBN 978-80-87342-04-6. ◾ MATELA, Jiří - RUSŇÁK, Vít - HOLUB, Petr. Efficient JPEG2000 EBCOT Context Modeling for Massively Parallel Architectures. In Storer, James A. and Marcellin, Michael W.. Data Compression Conference (DCC), 2011. Washington, DC, USA : IEEE Computer Society, 2011. od s. 423-432, 10 s. ISBN 978-0-7695-4352-9. ◾ MATELA, Jiří - ŠROM, Martin - HOLUB, Petr. Low GPU Occupancy Approach to Fast Arithmetic Coding in JPEG2000. In Z. Kotásek et al.. MEMICS 2011, LNCS 7119 - to appear. Heidelberg : Springer, 2011. od s. 136-145. 41/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Diskrétní vlnková transformace Rozložení obrazu do rekurzivně se opakujících pásů LL, HL, LH, HH signal in different resolutions • Most of advanced features of JPEG2000 rely on DWT • By application of 2D DWT, the source image is decomposed into four subbands (dentoted LL, HL, LH, HH) Result of application of 2D DWT 42/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Diskrétní vlnková transformace Lifting schéma ◾ low-pass a high-pass filtry d1 i = d0 i − 1 2 (s0 i + s0 i+1) s1 i = s0 i + 1 4 (d1 i−1 + d1 i ) Aplikace na řádky, poté na sloupce pass and high-pass filters d1 i = d0 i − 1 2(s0 i + s0 i+1) s1 i = s0 i + 1 4(d1 i−1 + d1 i ) ally filters are applied to each row ulting into coefficients of low-pass subband at even tions and coefficients of high-pass subband at odd tions → DWT Computation • Filters are applied to each column • Resulting into foursomes of coefficients of four subbands of DWT (LL, HL, LH, HH) 43/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Diskrétní vlnková transformace Lifting schéma ◾ přeuspořádání výsledného obrazu the global memory • Particular subbands however needs to be stored separately in global memory • It is important comply with global memory coalesced access • Even lines are stored first and first half of threads store LL coefficients and the second half stores HL coefficients Mapování na GPU ◾ 2D thread block ◾ každé vlákno zpracovává jeden lichý a jeden sudý prvek 44/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Diskrétní vlnková transformace Mapování na GPU that threads overlap only the upper half of the block, i.e., MAX(Ty) = Dy 2 − 1, which means that threads are directly mapped only to samples in the upper half of block data and we will have to change this mapping to be able to process both halves. Fig. 1. Source image partitioning. The first step of computation is to fetch image data from global memory into fast shared memory. It is crucial here to comply with coalesced global memory access. Considering the proposed data partitioning, each thread loads corresponding data sample into the upper half first, and then into the lower half of data block. The horizontal block size should be multiple of 16, so that coalesced access is not broken by thread block misalignment. DWT coefficients are then computed according to lifting scheme relations 1 and 2. To calculate first dimension of the transform, DWT filters are applied to every row separately. Afterwards, each row contains a sequence of interleaved coefficients of low-pass and high-pass subbands—L, H, L, H, ..., L, H, L, H. Each particular prediction and actualization step is calculated respectively as follows. 1. načtení z globální do sdílené paměti 2. aplikace lifting schéma na řádky Fig. 1. Source image partitioning. The first step of computation is to fetch image data from global memory into fast shared memory. It is crucial here to comply with coalesced global memory access. Considering the proposed data partitioning, each thread loads corresponding data sample into the upper half first, and then into the lower half of data block. The horizontal block size should be multiple of 16, so that coalesced access is not broken by thread block misalignment. DWT coefficients are then computed according to lifting scheme relations 1 and 2. To calculate first dimension of the transform, DWT filters are applied to every row separately. Afterwards, each row contains a sequence of interleaved coefficients of low-pass and high-pass subbands—L, H, L, H, ..., L, H, L, H. Each particular prediction and actualization step is calculated respectively as follows. s[Tx][2Ty + 1] = s[Tx][2Ty + 1] + p · (s[Tx][2Ty] + s[Tx][2Ty + 2]) (3) GPU-Based DWT Acceleration for JPEG2000 5 s[Tx][2Ty] = s[Tx][2Ty] + u · (s[Tx][2Ty − 1] + s[2Ty + 1]) (4) Where Tx and Ty determine the thread position in horizontal and vertical direction respectively and s[rowidx][columnidx] is the shared memory 2D array. Note that we propose transposed thread mapping for efficient data processing as follows. Threads are directly mapped into the upper half of block only, so that we have to change the thread mapping to be able to process whole block. In equations 3 and 4, we have swapped3 thread indices Tx, Ty so that the threads cover the left half of the data block instead of the upper half which was covered 3. aplikace lifting schéma na sloupce (prohození Tx a Ty) 4. uložení výsledků do globální paměti s přeuspořádáním uložení sudých pak lichých řádků v rámci řádku první polovina vláken čte LL (resp. LH), druhá polovina HL (resp. HH) ⇒ koalescentní přístup do globální paměti 45/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Diskrétní vlnková transformace• Performance of proposed GPU accelerated DWT compared to DWT in JasPer • JasPer is referential implementation of JPEG2000 • HD frame processed in 0.81 ms • About 68 times faster Implementation 512×512 1920×1080 Speedup JasPer 6ms 55ms N/A× CUDA DWT 0.12ms 0.81ms 67.9× 46/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Hledání kontextu pro kompresi aritmetickým adaptivním kodérem Jde se po jednotlivých bitplanech CM: Bit-Plane by Bit-Plane Source: wikipedia.org Bitplane skenován podle vzoru (scan-pattern) CM: Bit-Plane processed according to scan-pattern 47/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu 3 fáze: 1. Significance Propagation Pass (SPP) 2. Magnitude Refinement Pass (MRP) 3. Cleanup Pass (CUP) CM: Bit-Plane processed in three passes → CM: Bit-Plane processed in three passes → CM: Bit-Plane processed in three passes 48/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu 4 kódovací operace 1. Zero Coding (ZC) 2. Run-Length Coding (RLC) 3. Magnitude Refinement Coding (MRC) 4. Sign Coding (SC) každý bit je zakódován jednou nebo více operacemi právě v jedné fázi výstupem kódovacích operací je pár CX,D 49/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Určení fáze pro daný bit ◾ záleží na jeho stavu ◾ záleží na stavu jeho sousedů ◾ to vše se vyvíjí (proměnné σ, σ′ , η), jak bity prochází jednotlivými fázemi podle scan-pattern – :-((((((((((( Dle definice vysoce sekvenční proces ◾ umožňuje paralelizmus jedině na úrovni code block ◾ ... s jistými se speciálními výjimkami (causal mode) 50/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Reformulace problému ◾ náhrada původních stavových proměnných σ, σ′ , η za proměnné ρ, τ ◾ nové proměnné lze předpočítat paralelně pro každý bit ◾ výpočet ρ Computation of ρ ρp x,y = 1, if γx,y > 2p 0, otherwise Bitplane index 1 . . . n Pixel position in the code block Value of a pixel at x, y 51/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Reformulace problému ◾ výpočet τ Computation of τ • Fix-point algorithm simulating significance propagation of SPP 1. τp x,y = 1 ⇔ ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θA x,y (ρp+1 i,j = 1) 2. τp x,y = 1 ⇔ ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ4 x,y (τp+1 i,j = 1) 3. Step 2 is repeated until there is no new τp x,y = 1 found θA x,y θ4 x,y 52/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Reformulace problému ◾ pomocná proměnná δ Auxiliary state δ • δp x,y = 1 indicates a position γp x,y is in PN in SPP • δp x,y = (i,j)∈θA x,y (ρp+1 i,j = 1) ∨ (i,j)∈{θ5 x,y ,θ4 x,y ,θ3 x,y }(τp i,j = 1) • Where selection of surrounding θx,y depends on position of the bit γp x,y on y-axis. θA x,y θ5 x,y θ4 x,y θ3 x,y 53/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextu Důkazy ekvivalence ◾ Původní: Bit je kódován v SPP, pokud není signifikantní (σ = 0) a je v preferovaném okolí (preferred neighborhood). ◾ Paralelní: Bit je kódován v SPP, pokud není signifikantní (ρp+1 = 0) a je v preferovaném okolí (δp = 1). ◾ ...následuje důkaz ekvivalence indukcí ;-))) is processed in the four consecutive steps: MRC, RLC, ZC, SC. Note, that each coding operation is executed for all bits in a bit-plane in parallel. The only constraint on bit coding independence stems from diverging number of bits coded by the RLC operation. The RLC is defined to code one to four bits in column and a prediction of the number is virtually as expensive as the RLC coding itself. The only operation affected by this is ZC, so we choose to perform RLC operations in current bit-plane before ZC. Although the proposed design allows for parallelism among bit-planes too, we do not exploit it because of limited shared memory size on contemporary GPUs. The described fine-grained parallel algorithm allows for processing individual bits in parallel threads, resulting in high utilization of multi-processors on GPU. Depending on chosen code block size, the data may be processed entirely in the fast shared memory4 . V. Proof of Equivalence The equivalence of the original algorithm and the proposed parallel algorithm is based on equivalence of information provided by the original state variables σ, σ , and η and by the newly defined ρ, and τ . The proof is broken down into equivalency of states σ and ρ in already processed bit-planes (Lemma 1), and equivalency of σ and ρ ⊕ τ in SPP on current bit-plane (Lemma 2). The equivalence of information provided by η and ρ was shown in the description of MRP above. Replacement of information from σ by γP . . . γp+1 has also been described in MRP. Definition 5. A state variable σx,y is set to 1 (i.e., coefficient γx,y becomes significant) in the original algorithm right after the most significant non-zero bit γp x,y is processed in either SPP or CUP passes. A bit γp x,y becomes significant in SPP iff right before the bit is processed, it holds that σx,y = 0 ∧ γp x,y = 1 ∧ [ (i,j)∈θA x,y (σi,j = 1)]. Otherwise if σx,y = 0 ∧ γp x,y = 1, the bit becomes significant in CUP. Lemma 1. Value of σx,y is equal to ρp x,y prior to γp−1 x,y bit is processed. Proof. Value of σx,y is set to 1 for each coefficient γx,y once the most significant bit of the coefficient has been processed in the original sequential algorithm. I.e., right after processing bit γp x,y, σx,y = 1 iff one of bits γP−1 x,y . . . γp x,y was equal to 1, thus σx,y = P−1 p =p γp x,y. Following from Def. 1, ρp x,y = σx,y. In the sequential algorithm, this has to occur prior to proceeding to p − 1 bit-plane and thus also before γp−1 x,y is processed. Note, that σx,y of the original algorithm may change from 0 to 1 right after γp−1 x,y is processed and this is not reflected by ρp x,y, but it will be reflected by τp−1 x,y . Lemma 2. For a bit-plane p processed in the first pass (SPP) it holds that σx,y = 1 ⇔ ρp+1 x,y ⊕ τp x,y = 1, where ⊕ denotes XOR operation. Proof. Using mathematical induction we want to show that in each step of the first pass (SPP) in a bit-plane p: 4 Because of shared memory size limitations, older NVIDIA GPUs are limited to 16 × 16 code blocks, while new NVIDIA Fermi architecture allows for larger code blocks. 6 1. σx,y = 1 ⇒ ρp+1 x,y ⊕ τp x,y = 1 2. σx,y = 1 ⇐ ρp+1 x,y ⊕ τp x,y = 1 Note, that we assume that values of all state variables are equal to zero for indices x, y beyond borders of a bit-plane. Basis. Let x = 0, y = 0. 1. We want to show that if σ0,0 = 1 then ρp+1 0,0 ⊕ τp 0,0 = 1 A. Let σ0,0 = 1 before γp 0,0 is processed. Then σ0,0 = ρp+1 0,0 = 1 (Lemma 1). Following from Def. 3, τp 0,0 = 0 because ρp+1 0,0 = 1. B. Let σ0,0 = 1 right after, but not before, bit γp 0,0 is processed . Then coefficient γ0,0 becomes significant in current bit-plane p, and thus right before bit γp 0,0 is processed it holds that σ0,0 = 0 ∧ γp 0,0 = 1 ∧ [σ1,0 = 1 ∨ σ1,1 = 1 ∨ σ0,1 = 1]. Then following from Lemma 1, we can replace every σi,j with ρp+1 i,j for each unprocessed γp x,y hence ρp+1 0,0 = 0 ∧ γp 0,0 = 1 ∧ [ρp+1 1,0 = 1 ∨ ρp+1 1,1 = 1 ∨ ρp+1 0,1 = 1]. Therefore τp 0,0 = 1 from the first step of Definition 3. 2. We want to show that if ρp+1 0,0 ⊕ τp 0,0 = 1 then σ0,0 = 1 A. Let ρp+1 0,0 = 1 ∧ τp 0,0 = 0. Then σ0,0 = 1 from Definition 1 and Lemma 1. B. Let ρp+1 0,0 = 0 ∧ τp 0,0 = 1. Then coefficient γ0,0 becomes significant in current bit-plane p, and from Definition 3 we have ρp+1 0,0 = 0 ∧ γp 0,0 = 1 ∧ [ρp+1 1,0 = 1∨ρp+1 1,1 = 1∨ρp+1 0,1 = 1]. Then following from Lemma 1, we can replace every ρp+1 i,j with σi,j for each unprocessed γp x,y, hence, right before γp 0,0 is processed, it holds that σ0,0 = 0 ∧ γp 0,0 = 1 ∧ [σ1,0 = 1 ∨ σ1,1 = 1 ∨ σ0,1 = 1]. Thus σ0,0 = 1 right after γp 0,0 is processed (Def. 5). Induction step. We assume that σx ,y = 1 ⇔ ρp+1 x ,y ⊕ τp x ,y = 1 for each γp x ,y processed before γp x,y according to the prescribed scan pattern. 1. Let σx,y = 1. Then A. either coefficient γx,y became significant in one of previous bit-planes hence ρp+1 x,y = 1 (Lemma 1), B. or coefficient γx,y is going to become significant in the current bit-plane p, and therefore one of the three y mod 4 cases arises: I. ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ5 x,y (σi,j = 1) ∨ (i,j)∈θ3 x,y (ρp+1 i,j = 1) II. ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ4 x,y (σi,j = 1) ∨ (i,j)∈θ4 x,y (ρp+1 i,j = 1) III. ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ3 x,y (σi,j = 1) ∨ (i,j)∈θ5 x,y (ρp+1 i,j = 1) Using induction assumption we expand each σx,y to ρp+1 x,y ⊕ τp x,y = 1 so in case of 1.B.I we get ρp+1 x,y = 0∧γp x,y = 1∧ (i,j)∈θ5 x,y (ρp+1 i,j ⊕ τp i,j = 1) ∨ (i,j)∈θ3 x,y (ρp+1 i,j = 1) All terms ρp+1 i,j ⊕ τp i,j can be rewritten to a form of simple disjunction because both ρp+1 x,y and τp x,y cannot be equal to 1 at the same time. Hence ρp+1 x,y = 0 ∧ γp x,y = 1∧ (i,j)∈θ5 x,y (ρp+1 i,j = 1) ∨ (i,j)∈θ5 x,y (τp i,j = 1) ∨ (i,j)∈θ3 x,y (ρp+1 i,j = 1) 7 Both ρp+1 x,y can be combined together with respect to surroundings θ. Hence ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θA x,y (ρp+1 i,j = 1) ∨ (i,j)∈θ5 x,y (τp i,j = 1) Therefore τp x,y = 1 from Definition 3. The proof of the cases 1.B.II and 1.B.III is analogical to the 1.B.I. 2. Let ρp+1 x,y ⊕ τp x,y = 1. Then A. either ρp+1 x,y = 1, and thus coefficient γx,y became significant in one of previous bit-planes, hence σx,y = 1 (Lemma 1), B. or coefficient γx,y is going to become significant in current bit-plane, hence one of the three y mod 4 cases arises: I. ρp+1 x,y = 0 ∧ γp x,y = 1∧ (i,j)∈θ5 x,y (ρp+1 i,j ⊕ τp i,j = 1) ∨ (i,j)∈θ3 x,y (ρp+1 i,j = 1) II. ρp+1 x,y = 0 ∧ γp x,y = 1∧ (i,j)∈θ4 x,y (ρp+1 i,j ⊕ τp i,j = 1) ∨ (i,j)∈θ4 x,y (ρp+1 i,j = 1) III. ρp+1 x,y = 0 ∧ γp x,y = 1∧ (i,j)∈θ3 x,y (ρp+1 i,j ⊕ τp i,j = 1) ∨ (i,j)∈θ5 x,y (ρp+1 i,j = 1) Using induction assumption we expand each ρp+1 x,y ⊕ τp x,y = 1 to σx,y = 1 so in case of 2.B.I we get ρp+1 x,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ5 x,y (σi,j = 1) ∨ (i,j)∈θ3 x,y (ρp+1 i,j = 1) Each ρ can be rewritten to σ because those two state variables has same value before bit γp x,y is processed in current bit-plane p (Lemma 1). Hence σx,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θ5 x,y (σi,j = 1) ∨ (i,j)∈θ3 x,y (σi,j = 1) Surroundings θ5 x,y and θ3 x,y can be folded, hence σx,y = 0 ∧ γp x,y = 1 ∧ (i,j)∈θA x,y (σi,j = 1) Therefore following Def. 5 state σx,y = 1 right after bit γp x,y is coded. VI. Experimental Results Methodology. We set up two benchmark sets focused on the EBCOT Tier-1 processing speed only. Performance of bpcuda (our GPU implementation) was compared with two open-source CPU implementations (OpenJPEG5 1.2, JasPer6 1.900.1), one commercial (Kakadu7 2.2.3) and one GPU implementation (CUJ2K8 1.1). All the CPU implementations were single-threaded. 5 http://www.openjpeg.org 6 http://www.ece.uvic.ca/∼mdadams/jasper/ 7 http://www.kakadusoftware.com 8 http://cuj2k.sourceforge.net 8 54/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Modelování kontextuContext Modeling – Performance Color image in 1080p resolution OpenJPEG JasPer Kakadu CUJ2K CUDA GPU 0 200 400 600 800 642 ms, 107× 599 ms, 100× 98 ms 16× 94 ms 15× 6 ms, ref Performance[ms] 55/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Binární aritmetický kodér ◾ LPS do [0,Q) ◾ MPS do [Q,A] ◾ Q je pravděpodobnost výskytu LPS ◾ výsledek C Adaptivní ◾ mění se význam LPS/MPS Kontextový ◾ vstup má přiřazen kontext ◾ kontextu určuje stav kodéru mapování symbolů na MPS/LPS pravděpodobnost LPS MPS A = A(1 − Q) C = C + AQ LPS A = AQ C = C 56/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Problémy ◾ příliš hrubý paralelismus ◾ podmíněné větvené závislé na vstupních datech ◾ nedá se odstěhovat na CPU – na GPU máme tou dobou 2× tolik dat (CX,D) Pozitiva ◾ není tak náročné jako kontextové modelování MQ-Coder: Basic GPU Implementation • MQ-Coder implementation from OpenJPEG modified for GPU • Configuration: Every thread process the code-block • 1.9× faster then multi threaded CPU implementation 720p 1080p 4K Basic 38 ms 45 ms 93 ms (NVIDIA GeForce 580GTX)Nenašli jsme způsob principiální paralelizace :-( ◾ analýza různých optimalizačních technik ◾ někteří výzkumníci vymýšlejí jiné kodéry – problém nekompatibility se standardem 57/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Převod implementace do registrů ◾ lokální datové struktury do registrů ◾ průměrně: 240% zrychlení ⇒ registry jsou opravdu nejrychlejší (i vůči sdílené paměti) Rozvinování cyklů ◾ navrženo jako optimalizační technika pro VLSI ◾ zpracovává sérii MPS symbolů, pokud mají přiřazené stejné CX ◾ použití MAD instrukce: A = A − nQe, C = C + nQe ◾ paralelní prohledávání – detekce sekvencí ◾ průměrně: 31% zpomalení ⇒ zrychlení zpracování ani nevyváží režii předpočítání ⇒ zastoupení sérií MPS symbolů není dostatečně časté (17–25 % sekvencí 2–32) 58/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Prefix sum ◾ zpracovává sérii MPS symbolů nezávisle na CX ◾ nemůže používat MAD (různé CX znamená různé Qe) ◾ Qe jsou sčítány paralelně, součet použit pro úpravu A a C 59/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Renormalizace ◾ pokud A ≈ 1, můžeme aproximovat: A(1 − Q) ≈ A − Q a AQ = Q ◾ omezení počtu násobení Vylepšená rozšířená renormalizace ◾ renormalizace se provádí, pokud A klesne pod 0x8000 ◾ n-krát se násobí A a C, až je A > 0x8000 ◾ existující návrh urychlení: n určit vyhledávací tabulkou a použít shift ◾ publikovaný algoritmus měl chybu – neřešil korektně přetečení ◾ můžeme využít CLZ instrukci místo vyhledávací tabulky ◾ průměrně: 39% zrychlení ⇒ shifty a CLZ jsou rychlé 60/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Načítání dat po blocích ◾ hrubé vláknění vede na nekoalescentní přístup do paměti – vlákna zpracovávají různé části bloku ◾ pomůže načítat data po větších blocích1 ◾ využití double (8B) místo int (4B) ◾ optimalizace počtu natahovaných double podle výkonu ◾ průměrně: 33% zrychlení pro 16 natahovaných double 1Volkov, V.: Better Performance at Lower Occupancy. In: GPU Technology Conference 2010. (2010) 61/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Optimalizace využití karty ◾ díky hrubé granularitě každé vlákno potřebuje hodně zdrojů ◾ vytváří tlak na registry ◾ optimalizace využití karty na výsledný výkon 0 20 40 60 80 100 0.5 0.6 0.7 0.8 0.9 1 Occupancy [%] Relativeperformance 720p 1080p 4K 62/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Shrnutí optimalizací Low GPU Occupancy Approach to Fast Arithmetic Coding in JPE Fig. 4. Performance impact of individual optimizations (left) an performance relation (right). R – implementation in registers; ERN – rev renormalization; LU – loop unrolling; PS – prefix sum; CL – chunk load 720p 1080p 4K GPU Basic 38.0 ms – 45.0 ms – 92.9 ms – GPU R 16.2 ms 2.3× 18.9 ms 2.4× 48.5 ms 1.9× GPU R+ERN 11.9 ms 3.2× 14.9 ms 3.0× 44.9 ms 2.1× GPU R+LU 18.1 ms 2.1× 31.2 ms 1.4× 87.1 ms 1.1× GPU R+PS 20.9 ms 1.8× 25.7 ms 1.8× 64.6 ms 1.4× GPU R+CL 12.1 ms 3.1× 13.2 ms 3.4× 27.7 ms 3.4× GPU R+ERN+CL 7.3 ms 5.2× 8.1 ms 5.6× 17.6 ms 5.3× 0 20 0.5 0.6 0.7 0.8 0.9 1 Occ Relativeperformance 4 Conclusion and Future Work In this paper, we have studied various approaches to implementing a 63/63 Vlastnosti CUDA Metriky algoritmů JPEG2000 Aritmetické kódování – MQ-Coder Výsledný výkon • Combination of Register, Enhanced renormalization, and Chunk data loading • At 33% GPU occupancy • 5.6–16× speedup over multithreaded CPU implementation (Kakadu) • 3–9× speedup over GPU implementation (CUJ2K) 720p 1080p 4K OpenJPEG 1.4 157 ms 316 ms 1081 ms Jasper 1.900.1 89 ms 178 ms 594 ms Kakadu 6.4 (4 t) 41 ms 84 ms 284 ms CUJ2K 1.1 ≈25 ms ≈49 ms ≈166 ms CUDA GPU 7.3 ms 8.1 ms 17.6 ms