OpenCL - The open standard for parallel programming of heterogeneous systems
GPU-k általános számításokhoz GPU – Graphics Processing Unit Képalkotás: sok, általában egyszerű és független művelet < 2006: programmable shaders renderelés (Pixar RenderMan – Toy Story) 2006: CUDA -Compute Unified Device Architecture csak Nvidia kártyákon kezdetben 128 CUDA core, 8 compute unit (Tesla c870) 2008: OpenCL – cross-platform Khronos csoport (OpenGL!) OpenCL 1.0: Mac OS X Snow Leopard 2009 fontosabb partnerek: AMD, Intel (Larrabee) és természetesen Nvidia Alternatívák: DirectCompute, OpenMP, OpenACC stb.
OpenCL working group members Nvidia in chair, Apple is specification editor Már most is tekinthető általános szabványnak.
OpenCL könyvek OpenCL Programming Guide - The “Red Book” of OpenCL http://www.amazon.com/OpenCL-Programming-Guide-Aaftab-Munshi/dp/0321749642 OpenCL in Action http://www.amazon.com/OpenCL-Action-Accelerate-Graphics-Computations/dp/1617290173/ Heterogeneous Computing with OpenCL http://www.amazon.com/Heterogeneous-Computing-with-OpenCL-ebook/dp/B005JRHYUS The OpenCL Programming Book http://www.fixstars.com/en/opencl/book/
OpenCL API AMD http://developer.amd.com/zones/OpenCLZone/ INTEL http://software.intel.com/en-us/articles/opencl-sdk/ NVIDIA http://developer.nvidia.com/opencl
OpenCL: egy példa Soros megoldás
OpenCL megoldás
void add_vector(int dim,float *A,float *B, float *C) { for(int i=0;i
__kernel void add_vector(__global float *A, __global float *B,_global float *C) { int tidx = get_global_id(0);
Különbségek: ?
C[tidx]=A[tidx]+B[tidx]; }
OpenCL: egy példa (C) Soros megoldás
OpenCL megoldás
void add_vector(int dim,float *A,float *B, float *C) { for(int i=0;i
__kernel void add_vector(__global float *A, __global float *B,_global float *C) { int i = get_global_id(0); C[i]=A[i]+B[i]; }
Különbségek: - hol a vektor dimenziója? Soros kód: dim változó OpenCL: a get_global_id(0) maximumális értéke a dimenzió, melyet előzőleg beállítottunk a kernel hívás előtt - milyen memória címeket használhatunk? Minden OpenCL compute device saját memória területtel rendelkezik, melyet előre le kell foglalnunk
OpenCL: compute device memory types Global domain: work item size Local groups: workgroups Eszköz specifikus a paraméterezés, de mindent le tudunk kérdezni futásidőben.
OpenCL: memória Az architektúra legnagyobb hátránya: Egy számolás eredményének elérése: host → device → host Még abban az esetben is, ha a kódunkat a lokális CPU-n futtatjuk! A másolás egyik korlátja maga a PCI-express Lehetőség párhuzamos másolás-számolásra Megjegyzés: 1. sok esetben hatékonyabb helyben újraszámolni mint beolvasni 2. sok esetben a kiolvasás lassúsága miatt nem hatékony 3. Nincs virtuális memória → de lesz!
OpenCL: memória Hierarchikus memória: Minél Az architektúra legnagyobb hátránya:
lentebb megyünk annál kisebb a memória eredményének mérete, cserébe egyre Egy számolás elérése: gyorsabb! host → device → host Még 1. abban esetben is, ha a kódunkat Hostazmemory a lokálisCsak CPU-n futtatjuk! a CPU éri el, 10-20 GB/s
Akár 512 GB 2. Global/Constant memory Csak az adott device-on látható Lehetőség párhuzamos másolás-számolásra 100-250 GB/sec (GPU), <6GB 3. Local memory Megjegyzés: Csak egy adott workgroup látja 1. sok esetben hatékonyabb helyben 16-64 KB/work-item újraszámolni mint beolvasni 4. Private memory 2. sok esetben a kiolvasás lassúsága Csak az adott work-item látja miatt nemKB hatékony 16-64
A másolás egyik korlátja maga a PCI-express
3. Nincs virtuális memória → de lesz!
OpenCL: fordítás Fordítás történhet offline és online módon. Utóbbi az elterjedt. → univerzális OpenCL API: C/C++ → objektumok: Konfiguráció: Device - eszköz objektuma Context - eszközök → környezet Queue - feladatok kiosztása Memória: Buffer - memória blokkok Image - 2D vagy 3D kép Végrehajtás: Program - kernelek Kernel - maguk a feladatok
Példa fordításra: Mac OS X (>10.6): g++ -framework OpenCL cl_test.cpp -o cl_test Linux: Nvidia SDK g++ -I
/OpenCL/common/inc/ cl_test.cpp -o cl_test
OpenCL: fordítás
Példa: Alacsony szintű képi leírók számolása
Segmentation
Region of Interest
Dense Grid
Példa: Alacsony szintű képi leírók számolása Sok ezer független leíró képenként → a párhuzamosítás alapja Mivel GPU/CPU közös kód: - csak egy dimenziós párhuzamosítás - no image support Kis memória igény: - maga a nyers kép pl. 1Mio*3Byte - gradiens képek: 1Mio*4Byte - leírók: 128*4*N Byte Nyereség: ~10x CPU-hoz képest (nem optimális) Még több ha PCA + Fisher számolás : ~40x
