triSYCL is a research project to experiment with the specification of the SYCL standard and to give feedback to the Khronos Group SYCL_committee and also to the ISO C++ committee.
More recently, this project has been used to experiment with high-level C++ programming for AMD AIE CGRA inspired by some SYCL concepts, as explained in ACAP++ & AIE++: C++ extensions for AMD Versal AIE CGRA architecture.
Because of lack of resources this SYCL implementation is very incomplete and should not be used by a normal end-user. Fortunately there are now many other implementations of SYCL available, including some strong implementations like DPC++ or hipSYCL that can be used on various targets.
This implementation is mainly based on C++23 features backed with OpenMP or TBB for parallel execution on the CPU, with Boost.Compute for the non single-source OpenCL interoperability layer and with an experimental LLVM/Clang version for the device compiler (from 2017-2018 which is now obsolete) providing full single-source SYCL experience, typically targeting a 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.
Since around 2018 Intel has put a lot of effort in their own oneAPI DPC++ SYCL project to up-stream SYCL into LLVM/Clang, there is another project about merging the oneAPI DPC++ SYCL implementation with triSYCL at https://github.com/triSYCL/sycl to give a greater user experience for AMD FPGA instead of using our obsolete experimental clunky device compiler. But this is still very experimental because the AMD tool-chain is based on old incompatible versions of LLVM/Clang and nothing of these is supported by the AMD product teams.
triSYCL has been used to experiment and provide feedback for SYCL 1.2, 1.2.1, 2.2, 2020 and even the OpenCL C++ 1.0 kernel language from OpenCL 2.2.
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, then was mainly funded by Xilinx and now again by AMD since Xilinx has been acquired by AMD in 2022.
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. AMD is also hiring in this area... :-)
SYCL is a single-source modern C++-based DSEL (Domain Specific Embedded Language) and open standard from Khronos aimed at facilitating the programming of heterogeneous accelerators by leveraging existing concepts inspired by OpenCL, CUDA, C++AMP, OpenMP...
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.
SYCL is developed inside the Khronos SYCL committee and thus, for more information on SYCL, look at https://www.khronos.org/sycl
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 targeting various frameworks and API and has no relation with OpenCL itself, except when using the OpenCL API interoperability mode, like any other target.
For the SYCL ecosystem, look at https://sycl.tech
Most of our current efforts are focused on extensions, such as targeting AMD FPGA and Versal ACAP AIE CGRA, providing a way to program CPU, GPU, FPGA and CGRA at the same time in a single-source C++ program.
This project is a work-in-progress and currently we target partially only the first generation of devices, AIE/AIE1, while current models of AMD RyzenAI such as the Ryzen 9 7940HS has an AIE-ML/AIE2 as the XDNA/NPU/IPU.
The first generation programming model, ACAP++ was based on C++17/C++20 constructs.
See https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/test/acap/test_aie_mandelbrot.cpp and around, tests/acap and other directories starting with acap for some code samples running in pure C++ library CPU emulation with this project.
Look at doc/acap.rst to know more about how to install/use the ACAP++ environment.
The runtime for CPU emulation and AIE device is found in include/triSYCL/vendor/Xilinx which requires also a special compiler provided by https://github.com/triSYCL/sycl to run on VCK190 boards.
The second generation programming model, AIE++ is based on C++23/C++26 constructs, allowing an even terser syntax.
See around https://github.com/triSYCL/sycl/blob/sycl/unified/master/sycl/test/aie/mandelbrot.cpp for some examples.
The runtime for CPU emulation and AIE device support is found in include/aie and the compiler for device support is https://github.com/triSYCL/sycl
Other open-source projects related to AIE which are interesting to program AIE:
- https://riallto.ai
- https://github.com/Xilinx/mlir-aie
- https://github.com/Xilinx/mlir-air
- https://github.com/nod-ai/iree-amd-aie
Some documentation about AMD AIE CGRA:
- AIE aka AIE1
- Versal Adaptive SoC AI Engine Architecture Manual https://docs.xilinx.com/r/en-US/am009-versal-ai-engine/Overview
- Versal Adaptive SoC Technical Reference Manual https://docs.xilinx.com/r/en-US/am011-versal-acap-trm/Introduction
- AIE/AIE1 C++ API https://www.xilinx.com/htmldocs/xilinx2024_1/aiengine_api/aie_api/doc
- AIE-ML aka AIE2
- AIE2/AIE-ML architecture https://docs.xilinx.com/r/en-US/am020-versal-aie-ml/Overview
- AIE2/AIE-ML C++ API https://www.xilinx.com/htmldocs/xilinx2024_1/aiengine_ml_intrinsics/intrinsics
Please see about SYCL to have some context, a list of presentations, some related projects.
SYCL is a template library, so no real installation is required.
There are some examples you can build however.
See Testing.
Architecture of triSYCL runtime and compiler describes the code base with some high-level diagrams but also how it was possible to compile and use the obsolete device compiler on some AMD FPGA for example. Now look at https://github.com/triSYCL/sycl instead.
Some details about CMake configuration and organization can be found in CMake.
Yes, there are some macros used in triSYCL! Look at Pre-processor macros used in triSYCL to discover some of them.
See Environment variables with triSYCL.
See Possible futures.
The documentation of the triSYCL implementation itself can be found in https://trisycl.github.io/triSYCL/Doxygen/triSYCL/html and https://trisycl.github.io/triSYCL/Doxygen/triSYCL/triSYCL-implementation-refman.pdf
- 2023/06/09: merge the 5-year old branch experimenting with ACAP++ SYCL CPU model extensions for AMD Versal ACAP AIE1 CGRA like the XCVC1902 used in VCK190 or VCK5000 boards.
- 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 interest 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" (SPIR-df) and working at least with PoCL and Xilinx SDx xocc for FPGA.
- 2017/03/03: triSYCL can use CMake &
ctestand 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.