triSYCL

1   Introduction

triSYCL is an open source implementation to experiment with the specification of the OpenCL SYCL 1.2.1 and 2.2 C++ layer and to give feedback to the Khronos OpenCL SYCL and OpenCL C++ 2.2 kernel language committees and also ISO C++ committee.

This SYCL implementation is mainly based on C++17 and OpenMP for execution on the CPU, with Boost.Compute for the non single-source OpenCL interoperability layer and with LLVM/Clang for the device compiler providing single-source experience on SPIR device. Since in SYCL there is a host fall-back, this CPU implementation can be seen as an implementation of this fall-back too...

The parallel kernels can be executed in parallel on the CPU with OpenMP in the first range dimension, if compiled with OpenMP support or on an OpenCL device with the interoperability mode (which is not single source)

For legal reasons, the specification used for this open source project is the published current provisional specification and not the last one currently discussed in the Khronos OpenCL SYCL committee. If you are a Khronos member, you can access to https://gitlab.khronos.org/sycl/triSYCL where you might find more futuristic branches.

This is provided as is, without any warranty, with the same license as LLVM/Clang.

Technical lead: Ronan at keryell point FR. Developments started first at AMD and are now mainly funded by Xilinx.

It is possible to have a paid internship around triSYCL, if you have some skills related to this project. Contact the technical lead about this. Xilinx in also hiring in this area... :-)

2   OpenCL SYCL

OpenCL SYCL is a single-source C++14/C++17-based DSEL (Domain Specific Embedded Language) aimed at facilitating the programming of heterogeneous accelerators by leveraging the OpenCL language and concepts.

A typical kernel with its launch looks like this pure modern C++ code:

queue {}.submit([&](handler &h) {
    auto accA = bufA.get_access<access::mode::read>(h);
    auto accB = bufB.get_access<access::mode::write>(h);
    h.parallel_for<class myKernel>(myRange, [=](item i) {
        accA[i] = accB[i] + 1;
    });
});

Look for example at https://github.com/triSYCL/triSYCL/blob/master/tests/examples/demo_parallel_matrix_add.cpp for a complete example.

Note that even if the concepts behind SYCL are inspired by OpenCL concepts, the SYCL programming model is a very general asynchronous task graph model for heterogeneous computing with no relation with OpenCL itself, except when using the OpenCL API interoperability mode.

OpenCL SYCL is developed inside the Khronos OpenCL SYCL committee and thus, for more information on SYCL, look at http://www.khronos.org/sycl

For the SYCL ecosystem, look at http://sycl.tech

3   Documentation

3.1   Some reasons to use SYCL

Please see about SYCL

3.2   Installation & testing

SYCL is a template library, so no real installation is required.

There are some examples you can build however.

See Testing

3.3   Architecture of triSYCL runtime and compiler

Architecture of triSYCL runtime and compiler describes the code base with some high-level diagrams but also how to compile and use the device compiler on some Xilinx FPGA for example.

3.4   CMake infrastructure

Some details about CMake configuration and organization can be found in CMake

3.5   Pre-processor macros used in triSYCL

Yes, there are some macros used in triSYCL! Look at Pre-processor macros used in triSYCL to discover some of them.

3.6   Possible futures

See Possible futures

3.7   OpenCL triSYCL code documentation

The documentation of the triSYCL implementation itself can be found in http://xilinx.github.io/triSYCL/Doxygen/triSYCL/html and http://xilinx.github.io/triSYCL/Doxygen/triSYCL/triSYCL-implementation-refman.pdf

4   News

• 2018/03/12: the long-going device compiler branch has been merged in to provide experimental support for SPIR-df friendly devices, such as PoCL or Xilinx FPGA. This is only for the brave for now.
• 2018/02/01: there is now some documentation about the architecture of triSYCL on GPU and accelerators with its device compiler based on Clang/LLVM in doc/architecture.rst. While this is wildly experimental, there is a growing interested around it and it is always useful to get started as a contributor.
• 2018/01/05: there are some internship openings at Xilinx to work on triSYCL for FPGA https://xilinx.referrals.selectminds.com/jobs/compiler-engineer-intern-on-sycl-for-fpga-4685 and more generally Xilinx is hiring in compilation, runtime, C++, SYCL, OpenCL, machine-learning...
• 2017/12/06: the brand-new SYCL 1.2.1 specification is out and triSYCL starts moving to it
• 2017/11/17: the presentations and videos from SC17 on SYCL and triSYCL are now online https://www.khronos.org/news/events/supercomputing-2017
• 2017/09/19: there is a prototype of device compiler based on Clang/LLVM generating SPIR 2.0 "de facto" and working at least with PoCL and Xilinx SDx xocc for FPGA. Look at the device branch in https://github.com/triSYCL/triSYCL/tree/device
• 2017/03/03: triSYCL can use CMake & ctest and works on Windows 10 with Visual Studio 2017. It works also with Ubuntu WSL on Windows. :-) More info
• 2017/01/12: Add test case using the Xilinx compiler for FPGA
• 2016/11/18: If you missed the free SYCL T-shirt on the Khronos booth during SC16, you can always buy some on https://teespring.com/khronos-hpc (lady's sizes available, so no excuse! :-) )
• 2016/08/12: OpenCL kernels can be run with OpenCL kernel interoperability mode now.
• 2016/04/18: SYCL 2.2 provisional specification is out. This version implement SYCL 2.2 pipes and reservations plus the blocking pipe extension from Xilinx.