SYCL is an open industry standard, not a programming language itself, developed by the Khronos Group to facilitate data parallelism in C++ for heterogeneous computing systems. This means you can write single-source C++ code that utilizes the power of various processors like CPUs, GPUs, and FPGAs in a single program. This improves code maintainability and portability across different hardware platforms.

Core characteristics of SYCL:

  • Single-source programming: Write code once and target diverse hardware architectures without rewriting for each device.
  • Data parallelism: Leverage parallel execution for data-intensive tasks, significantly boosting performance.
  • Abstraction layer: SYCL hides the underlying hardware complexities, providing a unified programming model for different devices.
  • Memory management: SYCL offers various memory models (e.g., unified shared memory) to efficiently manage data movement between host and devices.
  • DPC++ (Data Parallel C++): Part of Intel’s oneAPI suite and also available as an open source compiler, DPC++ leverages LLVM/Clang for broad CPU and GPU support across Intel, NVIDIA, and AMD. In this post we will use DPC++ from oneAPI.
  • ComputeCPP: Open-source and multi-backend, ComputeCPP excels on Intel and AMD hardware with good CPU and GPU coverage. NVIDIA GPU support is under development.
  • hipSYCL: Primarily focused on NVIDIA GPUs with AMD support, hipSYCL offers multiple backends and experimental Intel GPU compatibility.

Additional technical points to consider

  • Performance portability: While SYCL promotes code portability, achieving optimal performance on each hardware platform might require specific tuning.
  • Learning curve: Compared to lower-level hardware-specific languages, SYCL has a gentler learning curve with support to modern C++ syntax but still requires understanding parallel programming concepts.
  • Ecosystem: The SYCL ecosystem is evolving, with growing adoption and toolchain improvements.
  • SYCL and IntelGPU in HPC: Aurora, the upcoming exascale supercomputer, will harness the power of SYCL and oneAPI on Intel GPUs to unlock groundbreaking scientific discoveries.

SYCL 2020

SYCL 2020 marked a significant milestone in the evolution of data parallelism for heterogeneous systems. DPC++ embraces the cutting-edge features introduced in the SYCL 2020 specification, empowering developers with:

  • Unified Shared Memory (USM): Streamlined memory management with seamless data access between host and devices.
  • Sub-groups: Fine-grained control over parallel execution within work groups for enhanced performance and algorithm optimization.
  • C++17 Syntax Simplifications (CTAD): Class template argument deduction for cleaner, more concise code.
  • Anonymous Lambdas: Direct use of lambda expressions without explicit naming, promoting code readability and flexibility.

These advancements collectively elevate the developer experience and performance capabilities of SYCL 2020, solidifying its position as a powerful and user-friendly tool for heterogeneous programming.

Key attributes of SYCL (DPC++)

  • Single source: Supports containing both host and device code in a single source file. We can also keep them in separate files, but it is not required.
  • Host: CPU. Need to have C++17 support to support DPC++ and SYCL programs.
  • Device: GPU/FPGA/DSP/ASIC/CPU. Used for acceleration offloading.
  • Sharing devices: As GPU device can be used by other programs, the performance of DPC++ program may be impacted.
  • Kernel code: Code that runs on accelerated devices.
  • List of features that are supported:
    • Lambdas
    • Operator overloading
    • Templates
    • Classes
    • Static polymorphism
    • Read only values of non-global host variables via lambda captures
  • The list of features not supported by kernel code are:
    • Dynamic polymorphism
    • Dynamic memory allocations (no new/delete operators)
    • Static variables
    • Function pointers
    • Runtime type information
    • Exception handling
    • No virtual member functions / variadic function can be called from kernel code
    • Recursion is not allowed within kernel code

Abstractions for Memory management in SYCL (DPC++)

  • Unified Shared Memory (USM)
    • Pointer based approach
    • Easier integration with existing C++ code that works on pointers
    • If we prefer to think about dependencies as performing one computation before another using an in-order queue, use USM
    • Allocation types:
    • device: Locates in device memory and cannot be accessed by host. Use malloc_device
    • host: Locates in host memory and accessible from both host and device. Use malloc_host
    • shared: Location shared and migrated automatically, accessible from both host and device. Use malloc_shared
    • Supports both explicit and implicit data movement based on allocation type
    • Explicit: Occurs in device allocation using handler’s memcpy() to transfer data from host to device and vice versa
    • Implicit: Occurs in host and shared allocations. No need to use memcpy() operation
  • Buffers
    • Represented by the buffer template class
    • Describe 1D/2D/3D arrays
    • Buffers follow implicit data movement strategy
    • If we prefer to think about data dependencies between kernels, choose buffers
    • Buffer object can be a scalar data type(int, float, double), vector+ data type or a user defined class or structure
    • Data structures in buffers must be C++ trivially copyable without the need of copy constructor
    • Can be accessed from host and device
    • Buffers represent data objects rather than specific memory addresses thus cannot be directly accessed like arrays
    • Buffers are not directly accessible, but can be accessed by accessor objects
    • Buffer access mode:
    • read: Read only access
    • write: Write only access. Previous content not disgarded
    • read_write: Read and write access
  • Images
    • Specific type of buffers for image processing
    • Supports image formats, reading of images using sampler objects

C++ Lambda Functions in SYCL

  • Lambda functions are introduced in C++11
  • Used to create anonymous function objects, can also be used with named variables
  • Lambda expression has the following syntax:

    [ capture-list ] ( params ) -> ret { body }
  • capture-list:

    • Lambda expression starts with a square bracket
    • Denotes how to capture variables that are used within the lambda but not passed as parameters
    • Comma separated list of captures
    • Variables from the surrounding scope mentioned in this list are available in the lambda body
    • Controls the visibility and lifetime of variables captured by lambda
    • Global variables are not captured in lambda expression
    • Non-global static variables can be used in a kernel but only if they are const
    • Capture a variable by value by listing the variable name in the capture-list. Even the value is modified in the body, it does not affect the original value outside the lambda.
    • Capture a variable by reference by listing the variable name prefixing with ampersand (&variable). If this variable is modified in the body, the original variable is also modified.
    • [=] shorthand to capture all variables used in the body by value and current object by reference
    • [&] shorthand to capture all variables used in the body by reference and current object by reference
    • [] captures nothing
    • SYCL uses [=] as kernel does not support capturing variable by reference
  • params:

    • List of function parameters similar to named function
    • In SYCL, it can be unique 1D id, or 2D/3D id
  • ret:

    • Defines the return type of the lambda expression
    • If -> ret is not specified, it is inferred from the return statement in the body
    • Return with no value implies void return type
    • SYCL kernels must always have a void return type thus it is not specified in SYCL kernel
  • body:

    • Contains function body
    • SYCL kernel body does not have any return statement

Use DPC++ using Docker in interactive mode

alt Use DPC++ using Docker in Interactive mode

Step 1: Pull Docker image

Pull Docker image for Intel oneAPI Basekit, which includes essential development tools. The process may take several minutes due to the image’s size (around a few gigabytes).:

docker pull intel/oneapi-basekit:devel-ubuntu22.04

Step 2: Prepare local environment

  • Open a terminal in a local directory where you plan to execute a sample SYCL program (e.g., vector_add.cpp).
  • Place your SYCL program in this directory. A sample program is provided below for reference:

    #include <sycl/sycl.hpp>
    using namespace sycl;
    void initialize_array(int n, double *a, double *b, double *c_result) {
    for (int i = 0; i < n; i++) {
        a[i] = 1.2;
        b[i] = 2.8;
        c_result[i] = a[i] + b[i];
    void vector_add(queue &Q, int n, double *a, double *b, double *c) {
    buffer<double, 1> a_buffer(a, range<1>(n));
    buffer<double, 1> b_buffer(b, range<1>(n));
    buffer<double, 1> c_buffer(c, range<1>(n));
    auto task_add = Q.submit([&](sycl::handler &cgh) {
        accessor a_accessor(a_buffer, cgh, read_only);
        accessor b_accessor(b_buffer, cgh, read_only);
        accessor c_accessor(c_buffer, cgh, write_only, no_init);
        cgh.parallel_for(range<1>(n), [=](id<1> idx) {
            c_accessor[idx] = a_accessor[idx] + b_accessor[idx];
    int main() {
    queue Q;
    std::cout << "Running on device: "
              << Q.get_device().get_info<info::device::name>()
              << std::endl;
    int n = 10000;
    double a[n];
    double b[n];
    double c[n];
    double c_result[n];
    initialize_array(n, a, b, c_result);
    vector_add(Q, n, a, b, c);
    int mismatch = 0;
    for (int i = 0; i < n; i++) {
        if (c_result[i] != c[i]) {
    if (mismatch == 0) {
        std::cout << "Verification passed" << std::endl;
    } else {
        std::cout << "Verification failed. Mismatch found: " << mismatch << std::endl;
    return 0;

Step 3: Mount the current directory to docker run

Mount the current directory into the Docker container to execute the SYCL program. The following command runs a container named inteldpc based on the intel/oneapi-basekit:devel-ubuntu22.04 image, mounting the current local directory to the /sycl_demo directory inside the container in interactive mode.

docker run -v $(pwd):/sycl_demo -it --name=inteldpc intel/oneapi-basekit:devel-ubuntu22.04

Step 4: List of available devices and compiler version in Docker system

  • Within the Docker container, list the available devices in the system using sycl-ls:

    root@5d3cac1a5b20:/# sycl-ls
    [opencl:acc:0] Intel(R) FPGA Emulation Platform for OpenCL(TM), Intel(R) FPGA Emulation Device 1.2 [2022.]
    [opencl:cpu:1] Intel(R) OpenCL, 13th Gen Intel(R) Core(TM) i9-13900H 3.0 [2022.]
  • Check the version of Intel LLVM-based compilers Intel® oneAPI DPC++/C++ Compiler:

    root@5d3cac1a5b20:/# icpx --version
    Intel(R) oneAPI DPC++/C++ Compiler 2023.0.0 (2023.0.0.20221201)
    Target: x86_64-unknown-linux-gnu
    Thread model: posix
    InstalledDir: /opt/intel/oneapi/compiler/2023.0.0/linux/bin-llvm
    Configuration file: /opt/intel/oneapi/compiler/2023.0.0/linux/bin-llvm/../bin/icpx.cfg

Step 5: Compile and execute SYCL program

  • Navigate to the mounted directory within the Docker container:

    root@5d3cac1a5b20:/# cd sycl_demo
  • Use the Intel oneAPI DPC++ compiler (icpx) to compile the SYCL program and run the executable:

    root@5d3cac1a5b20:/sycl_demo# icpx -fsycl vector_add.cpp
    root@5d3cac1a5b20:/sycl_demo# ./a.out
  • Upon successful execution, the program will display information about the running device and verification status:

    Running on device: 13th Gen Intel(R) Core(TM) i9-13900H
    Verification passed
  • If you exit the docker interactive mode and wants to turn on the interactive mode again you need to delete the container first:

    docker rm inteldpc

    Then follow the above steps again to compile and run SYCL program in docker interactive mode.

(Optional) Step 6: Compile and execute SYCL program using Makefile

  • It is easier to use the following Makefile to compile and run the program in the container interactive mode:

    CXX = icpx
    CXXFLAGS = -fsycl -std=c++17 -w
    SRCS := $(wildcard *.cpp)
    PRGS := $(patsubst %.cpp,%,$(SRCS))
    RUN :=  $(addprefix run_, $(PRGS))
    .PHONY: run_all
    run_all: $(sort $(RUN))
    .PHONY: all
    all: $(PRGS)
    .PHONY: all
    %: %.cpp
    	$(CXX) $(CXXFLAGS) -o $@ $^
    .PHONY: clean
    	rm -f $(PRGS)
  • Compile and run the program vector_add.cpp using make command:

    root@5d3cac1a5b20:/sycl_demo# make vector_add
    icpx -fsycl -std=c++17 -w -o vector_add vector_add.cpp
    root@5d3cac1a5b20:/sycl_demo# make run_vector_add
    Running on device: 13th Gen Intel(R) Core(TM) i9-13900H
    Verification passed

Some Useful Docker Commands

  • List and delete local Docker image:

    ➜  ~ docker image ls
    REPOSITORY                            TAG                                        IMAGE ID       CREATED        SIZE   7ed894ab0acc8ff09262113fdb08940d22654a30   82f0019ffd14   17 hours ago   13.4GB
    # Delete a specific image using its IMAGE ID
    ➜  ~ sudo docker image remove 82f0019ffd14 -f
    Deleted: sha256:82f0019ffd14bdfc269d5f01fc28eda497fc339d367448086f4441eeae634e7a
  • List and delete Docker container:

    # List all Docker containers, including stopped ones
    docker ps -a
    # Delete a specific container using its NAME or CONTAINER ID
    docker rm <NAME>
  • See disk usage in Docker:

    ➜  ~ docker system df
    TYPE            TOTAL     ACTIVE    SIZE      RECLAIMABLE
    Images          5         5         16.69GB   74.82MB (0%)
    Containers      7         0         20.95kB   20.95kB (100%)
    Local Volumes   2         2         50.18MB   0B (0%)
    Build Cache     0         0         0B        0B
  • Prune unused Docker data, including stopped containers, networks, and dangling images:

    ➜  ~ docker system prune
    WARNING! This will remove:
    - all stopped containers
    - all networks not used by at least one container
    - all dangling images
    - all dangling build cache
    Are you sure you want to continue? [y/N] y
    Deleted Containers:
    Deleted Networks:
    Total reclaimed space: 20.95kB




Click to select citation style

Shovon, A. R. (2023, December 26). Introduction to SYCL and DPC++. Ahmedur Rahman Shovon. Retrieved June 22, 2024, from

Shovon, Ahmedur Rahman. “Introduction to SYCL and DPC++.” Ahmedur Rahman Shovon, 26 Dec. 2023. Web. 22 Jun. 2024.

@misc{ shovon_2023,
    author = "Shovon, Ahmedur Rahman",
    title = "Introduction to SYCL and DPC++",
    year = "2023",
    url = "",
    note = "[Online; accessed 22-June-2024]"
Related contents in this website