Schneller Einstieg in OpenCL mit C++ Bindings

OpenCL Einfuhrung Ein ’einfaches’ OpenCL Programm OpenCL C++ Bindings OpenCL Bibliotheken

Schneller Einstieg in OpenCL mit C++ Bindings

Patrick CharrierBKW Consult

7. Mai 2014


Inhalt

1 OpenCL Einfuhrung

2 Ein ’einfaches’ OpenCL Programm

3 OpenCL C++ Bindings

4 OpenCL Bibliotheken


OpenCL Geschichte

Open Computing Language

Sprache fur Daten-parallele Programmierung

OpenCL 1.0 Spezifikation 2008

OpenCL 2.0 Spezifikation 2013

WebCL bringt OpenCL ins Web!

Unterstutzt

CPUs (AMD, Intel, ARM, . . . )GPUs (AMD, Intel, NVIDIA)Accelerators (Intel Xeon Phi, . . . )


OpenCL Platform-Modell

Host

Compute DeviceCompute UnitProcessing

ElementAbbildung: Copyright Advanced Micro Devices


OpenCL Kernels

Serieller Code

void s q u a r e ( i n t n , const f l o a t ∗a ,f l o a t ∗ r e s u l t )

{i n t i ;f o r ( i =0; i<n ; i ++)

r e s u l t [ i ] = a [ i ] ∗ a [ i ] ;}


OpenCL Kernels

Daten-paralleler Code mit OpenMP

void s q u a r e ( i n t n , const f l o a t ∗a ,f l o a t ∗ r e s u l t )

{#pragma omp p a r a l l e l f o rfo r ( i n t i =0; i<n ; i ++)

r e s u l t [ i ] = a [ i ] ∗ a [ i ] ;}


OpenCL Kernels

Daten-paralleler Code mit OpenCL

k e r n e l d p s q u a r e ( const f l o a t ∗a ,f l o a t ∗ r e s u l t )

{i n t i d = g e t g l o b a l i d ( 0 ) ;r e s u l t [ i d ] = a [ i d ] ∗ a [ i d ] ;

}

// dp squa r e e x e cu t e s oven ”n” work−i t ems


OpenCL Programmablauf

ContextContext

ProgramsPrograms KernelsKernels Memory*Objects

Memory*Objects

Command*Queue

Command*Queue

__kernel*void*sqr;__global*float*]input1*******__global*float*]output[{**size_t*id*=*get_global_id;0[;**output[id]*=*input[id]*]*input[id];}






sqrsqrarg[0]*valuearg[0]*valuearg[1]*valuearg[1]*value

imagesimagesimagesimagesimagesimages

imagesimagesimagesimagesbuffersbuffers

Compile Create*data*=*arguments Send*to*execution

In*Order*Queue

In*Order*Queue

Out*of*Order*Queue

Out*of*Order*Queue

Abbildung: Copyright Advanced Micro Devices


Ein ’einfaches’ OpenCL Programm

Arrays vertauschen in OpenCL

#inc l u d e <CL/ c l . h>

#de f i n e DATA SIZE 10

const char ∗ProgramSource =” k e r n e l v o i d swap ( g l o b a l f l o a t ∗a r ray1 , g l o b a l f l o a t ∗a r r a y 2 )\n”\”{\n”\” s i z e t i d = g e t g l o b a l i d (0 ) ;\n”\” f l o a t tmp = a r r a y2 [ i d ] ;\ n”\” a r r a y 2 [ i d ] = a r r a y 1 [ i d ] ;\ n”\” a r r a y 1 [ i d ] = tmp ;\n”\”}\n” ;



i n t main ( vo id ){c l c o n t e x t con t e x t ;c l c o n t e x t p r o p e r t i e s p r o p e r t i e s [ 3 ] ;c l k e r n e l k e r n e l ;c l command queue command queue ;c l p r og r am program ;c l i n t e r r ;c l u i n t num o f p l a t f o rms =0;c l p l a t f o r m i d p l a t f o rm i d ;c l d e v i c e i d d e v i c e i d ;c l u i n t num o f d e v i c e s =0;cl mem a1 , a2 ;s i z e t g l o b a l ;

f l o a t i nputData1 [ DATA SIZE]={1 , 2 , 3 , 4 , 5 , 6 , 7 , 8 , 9 , 10} ;f l o a t i nputData2 [ DATA SIZE ]={11 ,12 ,13 ,14 ,15 , 16 , 17 , 18 , 19 , 20} ;const s i z e t s i z e = DATA SIZE ∗ s i z e o f ( f l o a t ) ;

// r e t r e i v e s a l i s t o f p l a t f o rm s a v a i b l ei f ( c lG e tP l a t f o rm IDs (1 , &p l a t f o rm i d , &num o f p l a t f o rms ) != CL SUCCESS) {

p r i n t f ( ”Unable to ge t p l a t f o rm i d\n” ) ;r e t u r n 1 ;

}



// t r y to ge t a suppo r t ed GPU de v i c ei f ( c lGe tDev i c e ID s ( p l a t f o rm i d , CL DEVICE TYPE GPU , 1 , &d e v i c e i d , &

num o f d e v i c e s ) != CL SUCCESS) {p r i n t f ( ”Unable to ge t d e v i c e i d\n” ) ;r e t u r n 1 ;

}

// con t e x t p r o p e r t i e s l i s t − must be t e rm ina t ed wi th 0p r o p e r t i e s [0 ]= CL CONTEXT PLATFORM;p r o p e r t i e s [1 ]= ( c l c o n t e x t p r o p e r t i e s ) p l a t f o rm i d ;p r o p e r t i e s [2 ]= 0 ;

// c r e a t e a con t e x t w i th the GPU de v i c econ t e x t = c lC r e a t eCon t e x t ( p r o p e r t i e s ,1 ,& d e v i c e i d ,NULL ,NULL,& e r r ) ;

// c r e a t e command queue u s i n g the con t e x t and d e v i c ecommand queue = clCreateCommandQueue ( contex t , d e v i c e i d , 0 , &e r r ) ;

// c r e a t e a program from the k e r n e l s ou r c e codeprogram = clCreateProgramWithSource ( contex t , 1 , ( const char ∗∗) &ProgramSource ,

NULL , &e r r ) ;

// comp i l e the programi f ( c lBu i l dProg ram ( program , 0 , NULL , NULL , NULL , NULL) != CL SUCCESS) {

p r i n t f ( ” E r r o r b u i l d i n g program\n” ) ;r e t u r n 1 ;

}



// s p e c i f y which k e r n e l from the program to exe cu t ek e r n e l = c l C r e a t eK e r n e l ( program , ”swap” , &e r r ) ;

// c r e a t e b u f f e r s f o r the i n pu t and ouputa1 = c l C r e a t eB u f f e r ( contex t , CL MEM READ WRITE , s i z e , NULL , NULL) ;a2 = c l C r e a t eB u f f e r ( contex t , CL MEM READ WRITE , s i z e , NULL , NULL) ;

// l oad data i n t o the i n pu t b u f f e rc lEnqueueWr i t eBu f f e r ( command queue , a1 , CL TRUE , 0 , s i z e , inputData1 , 0 , NULL ,

NULL) ;c lEnqueueWr i t eBu f f e r ( command queue , a2 , CL TRUE , 0 , s i z e , inputData2 , 0 , NULL ,

NULL) ;

// s e t the argument l i s t f o r the k e r n e l commandc l S e tKe r n e lA r g ( k e r n e l , 0 , s i z e o f ( cl mem ) , &a1 ) ;c l S e tKe r n e lA r g ( k e r n e l , 1 , s i z e o f ( cl mem ) , &a2 ) ;g l o b a l=DATA SIZE ;

// enqueue the k e r n e l command f o r e x e c u t i o nclEnqueueNDRangeKernel ( command queue , k e r n e l , 1 , NULL , &g l oba l , NULL , 0 , NULL ,

NULL) ;c l F i n i s h ( command queue ) ;

// copy the r e s u l t s from out o f the output b u f f e rc lEnqueueReadBuf f e r ( command queue , a1 , CL TRUE , 0 , s i z e , inputData1 , 0 , NULL ,

NULL) ;c lEnqueueReadBuf f e r ( command queue , a2 , CL TRUE , 0 , s i z e , inputData2 , 0 , NULL ,

NULL) ;



// c l eanup − r e l e a s e OpenCL r e s o u r c e sc lRe leaseMemObject ( a1 ) ;c lRe leaseMemObject ( a2 ) ;c lRe l e a s eP rog ram ( program ) ;c l R e l e a s eK e r n e l ( k e r n e l ) ;clReleaseCommandQueue ( command queue ) ;c l R e l e a s eCon t e x t ( c on t e x t ) ;}

Das war

einfach

intuitiv

selbsterklarend

kurz

Oder doch nicht?



Dasselbe Programm in CUDA (ungetestet)

#de f i n e DATA SIZE 10

g l o b a l vo id swap ( g l o b a l f l o a t ∗a r ray1 , g l o b a l f l o a t a r r a y 2 ){

uns igned i n t i d = b l o c k I d x . x ∗ blockDim . x + th r e a d I d x . x ;f l o a t tmp = a r r a y2 [ i d ] ;a r r a y 2 [ i d ] = a r r a y 1 [ i d ] ;a r r a y 1 [ i d ] = tmp ;

}

(Kein Anspruch auf Richtigkeit)



i n t main ( ){

f l o a t i nputData1 [ DATA SIZE]={1 , 2 , 3 , 4 , 5 , 6 , 7 , 8 , 9 , 10} ;f l o a t i nputData2 [ DATA SIZE ]={11 ,12 ,13 ,14 ,15 , 16 , 17 , 18 , 19 , 20} ;const s i z e t s i z e = DATA SIZE ∗ s i z e o f ( f l o a t ) ;

// A l l o c a t e the d e v i c e i n pu t v e c t o r Af l o a t ∗d A = NULL ;cudaMal loc ( ( vo id ∗∗)&d A , s i z e ) ;

// A l l o c a t e the d e v i c e i n pu t v e c t o r Bf l o a t ∗d A = NULL ;cudaMal loc ( ( vo id ∗∗)&d B , s i z e ) ;

cudaMemcpy ( d A , inputData1 , s i z e , cudaMemcpyHostToDevice ) ;cudaMemcpy ( d B , inputData2 , s i z e , cudaMemcpyHostToDevice ) ;

i n t t h r e ad sPe rB l o ck = 256 ;i n t b l o c k sPe rG r i d = (DATA SIZE + th r ead sPe rB l o ck − 1) / th r e ad sPe rB l o ck ;swap<<<b l o ck sPe rG r i d , th r ead sPe rB lock>>>(d A , d B ) ;




cudaMemcpy ( inputData1 , d A , s i z e , cudaMemcpyDeviceToHost ) ;cudaMemcpy ( inputData2 , d B , s i z e , cudaMemcpyDeviceToHost ) ;

// Free d e v i c e g l o b a l memorycudaFree ( d A ) ;cudaFree ( d B ) ;

r e t u r n 0 ;}


Dem aufmerksamen Zuhorer/Leser fallt moglicherweise einUnterschied in der Lange des Codes auf.


OpenCL C++ Bindings

Analyse

CUDA ’versteckt’ mehr Details als OpenCL

PlatformDeviceContextCommand Queue

Das OpenCL Host-API ist ein C-API.OOP macht den Code kurzer.


OpenCL C++ Bindings

OpenCL Host-API Codes sind zu lang!

Moderne C++ Bindings schaffen Abhilfe:

OpenCL C++ Wrapper API (cl.hpp)openclamBoost.ComputeBOLT (AMD)

Alternative fur C-Programmierer: SimpleOpenCL


OpenCL C++ Wrapper API (cl.hpp)

Nicht standardisiert, aber auf Khronos-Website erhaltlich

Sehr nah am C-API, aber leichter lesbar

Ein wenig kurzer

Kein Performance-Overhead

//#d e f i n e NO STD VECTOR // Use c l : : v e c t o r i n s t e a d o f STL v e r s i o n#de f i n e CL ENABLE EXCEPTIONS#inc l u d e <CL/ c l . hpp>#inc l u d e <u t i l i t y >#inc l u d e <i o s t r eam>#inc l u d e <f s t r eam>#inc l u d e <s t r i n g>us ing namespace c l ;

i n t main ( ) {f l o a t i nputData1 [ DATA SIZE]={1 , 2 , 3 , 4 , 5 , 6 , 7 , 8 , 9 , 10} ;f l o a t i nputData2 [ DATA SIZE ]={11 ,12 ,13 ,14 ,15 , 16 , 17 , 18 , 19 , 20} ;const s i z e t s i z e = DATA SIZE ∗ s i z e o f ( f l o a t ) ;



t r y {// Get a v a i l a b l e p l a t f o rm svec to r<Plat form> p l a t f o rm s ;P la t fo rm : : ge t (& p l a t f o rm s ) ;

// S e l e c t the d e f a u l t p l a t f o rm and c r e a t e a con t e x t u s i n g t h i s p l a t f o rmand the GPU

c l c o n t e x t p r o p e r t i e s cps [ 3 ] = {CL CONTEXT PLATFORM,( c l c o n t e x t p r o p e r t i e s ) ( p l a t f o rm s [ 0 ] ) ( ) ,0

} ;Context c on t e x t ( CL DEVICE TYPE GPU , cps ) ;

// Get a l i s t o f d e v i c e s on t h i s p l a t f o rmvec to r<Device> d e v i c e s = con t e x t . g e t I n f o<CL CONTEXT DEVICES>() ;

// Crea te a command queue and use the f i r s t d e v i c eCommandQueue queue ( contex t , d e v i c e s [ 0 ] ) ;

// Read sou r c e f i l eProgram : : Sou rce s s ou r c e (1 , s t d : : make pa i r ( ProgramSource ,s t d : : s t r l e n ( ProgramSource )+1) ) ;



// Make program o f the s ou r c e code i n the con t e x tProgram program ( contex t , s ou r c e ) ;

// Bu i l d program f o r t h e s e s p e c i f i c d e v i c e sprogram . b u i l d ( d e v i c e s ) ;

// Make k e r n e lKerne l k e r n e l ( program , ”swap” ) ;

// Crea te memory b u f f e r sBu f f e r bu f f e rA ( contex t , CL MEM READ ONLY, s i z e ) ;Bu f f e r bu f f e rB ( contex t , CL MEM READ ONLY, s i z e ) ;

// Copy l i s t s A and B to the memory b u f f e r squeue . enqueueWr i t eBu f f e r ( bu f f e rA , CL TRUE , 0 , s i z e , bu f f e rA ) ;queue . enqueueWr i t eBu f f e r ( bu f f e rB , CL TRUE , 0 , s i z e , bu f f e rB ) ;

// Set arguments to k e r n e lk e r n e l . s e tArg (0 , bu f f e rA ) ;k e r n e l . s e tArg (1 , bu f f e rB ) ;

// Run the k e r n e l on s p e c i f i c ND rangeNDRange g l o b a l (DATA SIZE) ;NDRange l o c a l (1 ) ;queue . enqueueNDRangeKernel ( k e r n e l , Nul lRange , g l o ba l , l o c a l ) ;

// copy the r e s u l t s from out o f the output b u f f e rqueue . enqueueReadBuf fe r ( bu f f e rA , CL TRUE , 0 , s i z e , inputData1 ) ;queue . enqueueReadBuf fe r ( bu f f e rB , CL TRUE , 0 , s i z e , inputData2 ) ;



Exceptions, yeah!

} catch ( E r r o r e r r o r ) {s t d : : cout << e r r o r . what ( ) << ” ( ” << e r r o r . e r r ( ) << ” ) ” << s t d : : e nd l ;

}

r e t u r n 0 ;}


openclam

minimalistisch

STL kompatibel

kein zusatzlicher Compiler benotigt (vgl. CUDA)

#inc l u d e <openclam/ c l . hpp>#inc l u d e <vec to r>#inc l u d e <boos t / a s s i g n . hpp>

openclam : : openc l wrapper ;openclam : : c on t e x t con t e x t ( wrapper ) ;KERNEL( Add , contex t , f l o a t

k e r n e l vo id Add( g l o b a l const f l o a t∗ a ){

a [ g e t g l o b a l i d ( 0 ) ] += 42 ;} ; ) ;

s t d : : v e c to r< f l o a t > data = boos t : : a s s i g n : : l i s t o f ( 1 ) . r e p e a t ( 9 , 1 ) ;openclam : : f o r e a c h ( data . beg i n ( ) , data . end ( ) , Add ) ;


Boost.Compute

STL kompatibel

Implementiert die STL Algorithmen (count if, transform, etc.).

C++ Datenstrukturen konnen mit Macro in OpenCLverfugbar gemacht werden.

Unterstutzt Lambda Ausdrucke.

Events uber Boost.Thread Futures

Timing uber Boost.Chrono


Boost.Compute

#inc l u d e <vec to r>#inc l u d e <a l go r i t hm>#inc l u d e <boos t /compute . hpp>

namespace compute = boos t : : compute ;

i n t main ( ){

// get the d e f a u l t compute d e v i c ecompute : : d e v i c e gpu = compute : : system : : d e f a u l t d e v i c e ( ) ;

// c r e a t e a compute con t e x t and command queuecompute : : c on t e x t c t x ( gpu ) ;compute : : command queue queue ( ctx , gpu ) ;

// gen e r a t e random numbers on the hos ts t d : : v e c to r<f l o a t> h o s t v e c t o r (1000000) ;s t d : : g en e r a t e ( h o s t v e c t o r . b eg i n ( ) , h o s t v e c t o r . end ( ) , rand ) ;

// c r e a t e v e c t o r on the d e v i c ecompute : : v e c to r<f l o a t> d e v i c e v e c t o r (1000000 , c t x ) ;


Boost.Compute

// copy data to the d e v i c ecompute : : copy (

h o s t v e c t o r . b eg i n ( ) ,h o s t v e c t o r . end ( ) ,d e v i c e v e c t o r . b eg i n ( ) ,queue

) ;

// s o r t data on the d e v i c ecompute : : s o r t (

d e v i c e v e c t o r . b eg i n ( ) ,d e v i c e v e c t o r . end ( ) ,queue

) ;

// copy data back to the hos tcompute : : copy (

d e v i c e v e c t o r . b eg i n ( ) ,d e v i c e v e c t o r . end ( ) ,h o s t v e c t o r . b eg i n ( ) ,queue

) ;

r e t u r n 0 ;}


Mathematische OpenCL Bibliotheken

Mehrere OpenCL Bibliothekenhaben integrierte C++ Bindings.

Hauptzweck ist Parallelisierungrechenaufwendiger mathematischer Aufgaben.

Nachteil von Bibliotheken: Großer als Bindings


Mathematische OpenCL Bibliotheken

Allrounder:

ViennaCLVexCL

Spezialisten:

COGL (Sorting, Scanning)MAGMA (Matrix-Operationen)RajinCL (Matrix-Operationen)libCL (Filter, SPH, Radix, BVH, Vektoren)OpenCLIPP (Bildverarbeitung)ArrayFire (Array-Operationen)


ViennaCL

Hat Schnittstellen zuSTL, Eigen, Boost.uBLAS, MTL, MATLAB, Python

Alternative zu cublas, cufft, etc.

Unterstutzt

BLAS OperationenDichte and dunnbesetzte MatrizenDirekte und iterative LoserSingularwertzerlegungQR FaktorisierungFast Fourier Transform


VexCL

Sehr machtig! Das ’Eigen’ fur OpenCL.

Nutzt Expression Templates (wie Eigen).

Hat Schnittstellen zuSTL, Eigen, Boost.Compute, ViennaCL, CLOGS

Unterstutzt:

’Lazy’ Matrix Slicing, ReshapingDunnbesetzte Matrix-Vektor Produkte’Multivektoren’(Vektor von Vektoren fur kompakte Aufrufe von Kerneln)Reduction, Sorting, ScanningB-SplinesFast Fourier TransformKernel Generator (Beispiel: Runge Kutta hohe Ordnung)


Vielen Dank!

Fragen?


OpenCL Speicher-Modell

Host Memory

Global/Constant Memory

Local Memory

Work-item

Work-item

Private Memory

Private Memory

Workgroup

Host

Compute Device

Local Memory

Work-item

Work-item

Private Memory

Private Memory

Workgroup

Abbildung: Copyright Advanced Micro Devices

Schneller Einstieg in OpenCL mit C++ Bindings

Software

Transcript of Schneller Einstieg in OpenCL mit C++ Bindings