WIXLIB

C dialects for OpenCL (and heterogeneous computing)IC AMPparallel for each (e ,/ / D e f i n e t h e k e r n e l t o execute[ ] ( Concurrency : : index 1 i ) r e s t r i c t (amp) {a[ i ] i [0];});/ / I n t h e d e s t r u c t i o n o f a r r a y v i e w " a " happening here ,/ / t h e data are copied back b e f o r e i o t a n ( ) r e t u r n s/ / Use i o t a a l g o r i t h m i n C AMP# i n c l u d e amp . h # i n c l u d e iostream enum { NWITEMS 512 } ;i n t data [NWITEMS ] ;}/ / To a v o i d w r i t i n g Concurrency : : everywhereu s i n g namespace Concurrency ;void iota n ( size t n , i n t dst [ ] ) {/ / S e l e c t t h e f i r s t t r u e a c c e l e r a t o r found as t h e d e f a u l t onef o r ( auto c o n s t & acc : a c c e l e r a t o r : : g e t a l l ( ) )i f ( ! acc . g e t i s e m u l a t e d ( ) ) {a c c e l e r a t o r : : s e t d e f a u l t ( acc . g e t d e v i c e p a t h ( ) ) ;break ;}/ / D e f i n e t h e i t e r a t i o n spacee x t e n t 1 e ( n ) ;/ / Create a b u f f e r from t h e g i v e n a r r a y memoryarray view i n t , 1 a ( e , d s t ) ;/ / I s t h e r e a b e t t e r way t o express w r i t e o n l y data ?a . discard data ( ) ;/ / Execute a k e r n e l i n p a r a l l e l Modern C , heterogeneous computing & OpenCL SYCL Developed by Microsoft, AMD &MultiCoreWare Single source: easy to write kernels Require specific compiler Not pure C (restrict, tile static) No OpenCL interoperability Difficult to optimize the data transfersIWOCL 201521 / 43

C dialects for OpenCL (and heterogeneous computing)IOpenMP 4# i n c l u d e s t d i o . h enum { NWITEMS 512 } ;i n t a r r a y [NWITEMS ] ;p r i n t f ( "%d %d \ n " , i , a r r a y [ i ] ) ;return 0;}void iota n ( size t n , i n t dst [ n ] ) { Old HPC standard from the 90’s#pragma omp t a r g e t map( from : d s t [ 0 : n 1]) Use #pragma to express parallelism#pragma omp p a r a l l e l f o r OpenMP 4 extends it to acceleratorsf o r ( i n t i 0 ; i n ; i )dst [ i ] i ;I Work-group parallelism}I Work-item parallelismi n t main ( i n t argc , c o n s t char argv [ ] ) { Deal with CPU & heterogeneouscomputing parallelismi o t a n (NWITEMS, a r r a y ) ;/ / Display r e s u l t sf o r ( i n t i 0 ; i NWITEMS; i ) Modern C , heterogeneous computing & OpenCL SYCL No LDS support No OpenCL interoperability But quite simple! Single source.IWOCL 201522 / 43

C dialects for OpenCL (and heterogeneous computing)IOther (non-)OpenCL C framework(I) ArrayFire, Aura, CLOGS, hemi, HPL, Kokkos, MTL4, SkelCL, SkePU, EasyCL. nVidia CUDA 7 now C 11-based;I Single sourcesimpler for the programmerI nVidia Thrust parallel STL map-reduce on top of CUDA, OpenMP or TBBhttps://github.com/thrust/thrust Not very clean because device pointers returned by cudaMalloc() do not have a special typeuse some ugly casts; OpenACC OpenMP 4 restricted to accelerators LDS finer control Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201523 / 43

C dialects for OpenCL (and heterogeneous computing)IMissing link. No tool providingI OpenCL interoperabilityI Modern C environmentI Single source for programming productivity Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201524 / 43

OpenCL SYCL 1.2IOutline1C 142C dialects for OpenCL (and heterogeneous computing)3OpenCL SYCL 1.2C . putting everything altogether4OpenCL SYCL 2.1.5Conclusion Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201525 / 43

OpenCL SYCL 1.2IPuns and pronunciation explained Modern C , heterogeneous computing & OpenCL SYCLOpenCL SYCLOpenCL SPIRsickle [ "si-k@l ]spear [ "spir ]IWOCL 201526 / 43

OpenCL SYCL 1.2IOpenCL SYCL goals Ease of useI Single source programming model Take advantage of CUDA & C AMP simplicity and power Compiled for host and device(s) Easy development/debugging on host: host fall-back target Programming interface based on abstraction of OpenCL components (datamanagement, error handling.) Most modern C features available for OpenCLI Enabling the creation of higher level programming modelsI C templated libraries based on OpenCLI Exceptions for error handling Portability across platforms and compilers Providing the full OpenCL feature set and seamless integration with existing OpenCLcode Task graph programming model with interface à la TBB/Cilk (C 17) High performancehttp://www.khronos.org/opencl/sycl Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201527 / 43

OpenCL SYCL 1.2IComplete example of matrix addition in OpenCL SYCL# i n c l u d e CL / s y c l . hpp # i n c l u d e iostream b u f f e r f l o a t , 2 B { b , range 2 { N, M } } ;b u f f e r f l o a t , 2 C { c , range 2 { N, M } } ;/ / Enqueue some computation k e r n e l t a s kmyQueue . submit ( [ & ] ( h a n d l e r& cgh ) {/ / D e f i n e t h e data used / producedauto ka A . get access access : : read ( cgh ) ;auto kb B . get access access : : read ( cgh ) ;auto kc C . get access access : : w r i t e ( cgh ) ;/ / Create & c a l l OpenCL k e r n e l named " mat add "cgh . p a r a l l e l f o r c l a s s mat add ( range 2 { N, M } ,[ ] ( i d 2 i ) { kc [ i ] ka [ i ] kb [ i ] ; });} ) ; / / End o f our commands f o r t h i s queue} / / End scope , so w a i t f o r t h e queue t o complete ./ / Copy back t h e b u f f e r data w i t h RAII b e h a v i o u r .return 0;u s i n g namespace c l : : s y c l ;constexpr s i z e t N 2;constexpr s i z e t M 3;u s i n g M a t r i x f l o a t [N ] [M] ;i n t main ( ) {Matrix a { { 1 , 2 , 3 } , { 4 , 5 , 6 } } ;Matrix b { { 2 , 3 , 4 } , { 5 , 6 , 7 } } ;Matrix c ;{ / / Create a queue t o work onqueue myQueue ;/ / Wrap some b u f f e r s around our datab u f f e r f l o a t , 2 A { a , range 2 { N, M } } ; Modern C , heterogeneous computing & OpenCL SYCL}IWOCL 201528 / 43

OpenCL SYCL 1.2IAsynchronous task graph model Theoretical graph of an application described implicitly with kernel tasks using buffersthrough accessorscl::sycl::accessor write init acl::sycl::accessor read cl::sycl::accessor write cl::sycl::buffer ainit bcl::sycl::buffer bcl::sycl::buffer cmatrix addcl::sycl::accessor write Displaycl::sycl::accessor read cl::sycl::accessor read Possible schedule by SYCL runtime:init b init a matrix addDisplay; Automatic overlap of kernels & communicationsI Even better when looping around in an applicationI Assume it will be translated into pure OpenCL event graphI Runtime uses as many threads & OpenCL queues as necessary (AMD synchronousqueues, AMD compute rings, AMD DMA rings.) Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201529 / 43

OpenCL SYCL 1.2ITask graph programming — the code# i n c l u d e CL / s y c l . hpp # i n c l u d e iostream u s i n g namespace c l : : s y c l ;/ / Size o f t h e m a t r i c e sc o n s t s i z e t N 2000;c o n s t s i z e t M 3000;i n t main ( ) {{ / / By s t i c k i n g a l l t h e SYCL work i n a { } block , we ensure/ / a l l SYCL t a s k s must complete b e f o r e e x i t i n g t h e b l o c k[ ] ( auto i n d e x ) {B [ i n d e x ] i n d e x [0] 2014 i n d e x [ 1 ] 4 2 ;});});/ / Launch an asynchronous k e r n e l t o compute m a t r i x a d d i t i o n c a bmyQueue . submit ( [ & ] ( auto &cgh ) {/ / I n t h e k e r n e l a and b are read , b u t c i s w r i t t e nauto A a . get access access : : read ( cgh ) ;auto B b . get access access : : read ( cgh ) ;auto C c . get access access : : w r i t e ( cgh ) ;/ / From these accessors , t h e SYCL r u n t i m e w i l l ensure t h a t when/ / t h i s k e r n e l i s run , t h e k e r n e l s computing a and b completed/ / Create a queue t o work onqueue myQueue ;/ / Create some 2D b u f f e r s o f f l o a t f o r our m a t r i c e sb u f f e r double , 2 a ( { N, M } ) ;b u f f e r double , 2 b ( { N, M } ) ;b u f f e r double , 2 c ( { N, M } ) ;/ / Launch a f i r s t asynchronous k e r n e l t o i n i t i a l i z e amyQueue . submit ( [ & ] ( auto &cgh ) {/ / The k e r n e l w r i t e a , so g e t a w r i t e accessor on i tauto A a . get access access : : w r i t e ( cgh ) ;/ / Enqueue a p a r a l l e l k e r n e l on a N M 2D i t e r a t i o n spacecgh . p a r a l l e l f o r c l a s s matrix add ( { N, M } ,[ ] ( auto i n d e x ) {C[ index ] A[ index ] B[ index ] ;});});/ Request an access t o read c from t h e host s i d e . The SYCL r u n t i m eensures t h a t c i s ready when t h e accessor i s r e t u r n e d /auto C c . get access access : : read , access : : h o s t b u f f e r ( ) ;s t d : : c o u t s t d : : e n d l " R e s u l t : " s t d : : e n d l ;f o r ( s i z e t i 0 ; i N ; i )f o r ( s i z e t j 0 ; j M; j )/ / Compare t h e r e s u l t t o t h e a n a l y t i c v a l u ei f (C [ i ] [ j ] ! i (2 2014) j (1 4 2 ) ) {s t d : : c o u t " Wrong v a l u e " C [ i ] [ j ] " on element " i ’ ’ j s t d : : e n d l ;e x i t ( 1);}} / End scope o f myQueue , t h i s w a i t f o r any r e m a i n i n g o p e r a t i o n s on t h equeue t o complete /s t d : : c o u t " Good computation ! " s t d : : e n d l ;return 0;/ / Enqueue p a r a l l e l k e r n e l on a N M 2D i t e r a t i o n spacecgh . p a r a l l e l f o r c l a s s i n i t a ( { N, M } ,[ ] ( auto i n d e x ) {A [ i n d e x ] i n d e x [0 ] 2 i n d e x [ 1 ] ;});});/ / Launch an asynchronous k e r n e l t o i n i t i a l i z e bmyQueue . submit ( [ & ] ( auto &cgh ) {/ / The k e r n e l w r i t e b , so g e t a w r i t e accessor on i tauto B b . get access access : : w r i t e ( cgh ) ;/ From t h e access p a t t e r n above , t h e SYCL r u n t i m e d e t e c tt h i s command group i s independant from t h e f i r s t oneand can be scheduled i n d e p e n d e n t l y // / Enqueue a p a r a l l e l k e r n e l on a N M 2D i t e r a t i o n spacecgh . p a r a l l e l f o r c l a s s i n i t b ( { N, M } , Modern C , heterogeneous computing & OpenCL SYCL}IWOCL 201530 / 43

OpenCL SYCL 1.2IFrom work-groups & work-items to hierarchical parallelismconst i n t size 10;i n t data [ s i z e ] ;const i n t gsize 2;b u f f e r i n t my buffer { data , s i z e } ;Very close to OpenMP 4 style! ,my queue . submit ( [ & ] ( auto &cgh ) {auto i n my buffer . get access access : : read ( cgh ) ;auto o u t my buffer . get access access : : w r i t e ( cgh ) ;/ / I t e r a t e on t h e work groupcgh . p a r a l l e l f o r w o r k g r o u p c l a s s h i e r a r c h i c a l ( { s i z e ,gsize } ,[ ] ( group grp ) {/ / Code executed o n l y once per work groups t d : : c e r r " Gid " grp [ 0 ] s t d : : e n d l ;/ / I t e r a t e on t h e work i t e m s o f a work groupcgh . p a r a l l e l f o r w o r k i t e m ( grp , [ ] ( item 1 t i l e ) {s t d : : c e r r " i d " t i l e . g e t l o c a l ( ) [ 0 ] " " t i l e . g e t g l o b a l ( ) [ 0 ] s t d : : e n d l ;out [ t i l e ] i n [ t i l e ] 2;});/ / Can have o t h e r cgh . p a r a l l e l f o r w o r k i t e m ( ) here . . .});}); Modern C , heterogeneous computing & OpenCL SYCL Easy to understand the concept ofwork-groups Easy to write work-group only code Replace code barriers withseveral parallel for workitem ()I Performance-portable betweenCPU and GPUI No need to think about barriers(automatically deduced)I Easier to compose components &algorithmsI Ready for future GPU with nonuniform work-group sizeIWOCL 201531 / 43

OpenCL SYCL 1.2IC 11 allocators C 11 allocators to control the way objects are allocated in memoryI For example to allocate some vectors on some storageI Concept of scoped allocator to control storage of nested data structuresI Example: vector of strings, with vector data and string data allocated in different memoryareas (speed, power consumption, caching, read-only.) SYCL reuses allocator to specify how buffer and image are allocated on the hostside Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201532 / 43

OpenCL SYCL 1.2IC . putting everything altogetherOutline1C 142C dialects for OpenCL (and heterogeneous computing)3OpenCL SYCL 1.2C . putting everything altogether4OpenCL SYCL 2.1.5Conclusion Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201533 / 43

OpenCL SYCL 1.2IC . putting everything altogetherExascale-ready Use your own C compilerI Only kernel outlining needs SYCL compiler SYCL with C can address most of the hierarchy levelsI MPII OpenMPI C -based PGAS (Partitioned Global Address Space) DSeL (Domain-Specific embeddedLanguage, such as Coarray C .)I Remote accelerators in clustersI Use SYCL buffer allocator for RDMAOut-of-core, mapping to a filePiM (Processor in Memory). Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201534 / 43

OpenCL SYCL 1.2IC . putting everything altogetherDebugging Difficult to debug code or detect precondition violation on GPU and at large. Rely on C to help debuggingI Overload some operations and functions to verify preconditionsI Hide tracing/verification code in constructors/destructorsI Can use pure-C host implementation for bug-tracking with favorite debugger Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201535 / 43

OpenCL SYCL 1.2IC . putting everything altogetherPoor-man SVM with C 11 SYCL For complex data structuresI Objects need to be in buffers to be shipped between CPU and devicesI Do not want marshaling/unmarshaling objects.I Use C 11 allocator to allocate some objects in 1 SYCL buffer Useful to send efficiently data through MPI and RDMA too!I But since no SVM, not same address on CPU and GPU side. How to deal with pointers? / Override all pointer accessed (for example use std::pointer trait) to do address translationon kernel side ,Cost: 1 addition per *p When no or inefficient SVM.I Also useful optimization when need to work on a copy only on the GPU Only allocation on GPU side Spare some TLB trashing on the CPU Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201536 / 43

OpenCL SYCL 1.2IC . putting everything altogether¿¿¿Fortran? Fortran 2003 introduces C-interoperability that can be used for C interoperability.SYCL C boost:: multi array & others provides à la Fortran arraysI Allows triplet notationI Can be used from inside SYCL to deal with Fortran-like arrays Perhaps the right time to switch your application to modern C ? , Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201537 / 43

OpenCL SYCL 1.2IC . putting everything altogetherUsing SYCL-like models in other areas SYCL generic heterogeneous computing model beyond OpenCLIIIIIqueue expresses where computations happenparallel for launches computationsaccessor defines the way we access databuffer for storing dataallocator for defining how data are allocated/backed Example for HSA: almost direct mapping à la OpenCL Example in PiM worldI Use queue to run on some PiM chipsI Use allocator to distribute data structures or to allocate buffer in special memory(memory page, chip.)I Use accessor to use alternative data access (split address from computation, streamingonly, PGAS.)I Use pointer trait to use specific way to interact with memoryI . Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201538 / 43

OpenCL SYCL 2.1.IOutline1C 142C dialects for OpenCL (and heterogeneous computing)3OpenCL SYCL 1.2C . putting everything altogether4OpenCL SYCL 2.1.5Conclusion Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201539 / 43

OpenCL SYCL 2.1.ISYCL 2.1 is coming! Skip directly to OpenCL 2.1 and C 14 Kernel side enqueue Shared memory between host and accelerator Parallel STL C 17 Array TS Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201540 / 43

OpenCL SYCL 2.1.ISYCL and fine-grain system shared memory (OpenCL 2)# i n c l u d e CL / s y c l . hpp # i n c l u d e iostream # include vector u s i n g namespace c l : : s y c l ;i n t main ( ) {std : : vector a { 1 , 2 , 3 } ;std : : vector b { 5 , 6 , 8 } ;std : : vector c ( a . size ( ) ) ;/ / Enqueue a p a r a l l e l k e r n e lp a r a l l e l f o r ( a . size ( ) , [ & ] ( i n t index ) {c [ index ] a [ index ] b [ index ] ;});/ / Since t h e r e i s no queue o r no accessor , we assume p a r a l l e l f o r are b l o c k i n g k e r n e l ss t d : : c o u t s t d : : e n d l " R e s u l t : " s t d : : e n d l ;f o r ( auto e : c )s t d : : c o u t e " " ;s t d : : c o u t s t d : : e n d l ;return 0;} Very close to OpenMP simplicity Can still use of buffers & accessors for compatibility & finer control (task graph,optimizations.)I SYCL can remove the copy when possible Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201541 / 43

ConclusionIOutline1C 142C dialects for OpenCL (and heterogeneous computing)3OpenCL SYCL 1.2C . putting everything altogether4OpenCL SYCL 2.1.5Conclusion Modern C , heterogeneous computing & OpenCL SYCLIWOCL 201542 / 43

ConclusionIConclusion Many C frameworks to leverage OpenCLI None of them provides seamless single source Require some kind of macros & weird syntaxI But they should be preferred to plain OpenCL C for productivity SYCL provides seamless single source with OpenCL

IWOCL 2015 SYCL Tutorial C 14 I Outline 1 C 14 2 C dialects for OpenCL (and heterogeneous computing) 3 OpenCL SYCL 1.2 C . putting everything altogether 4 OpenCL . IParallelism TS (Technical Specification) with Parallel STL IConcurrency TS (threads, mutex.) IArray TS (multidimensional arrays à la Fortran) ITransactional Memory TS .