Introduction to SYCL

Device Memory using Buffers and Accessors

Setup Device Memory - Buffers


In order to handle data on the device memory we need to create SYCL buffers.

The constructor for buffers for a raw data pointer to the data to be managed is the following:

cl::sycl::buffer<typename T, int dimensions>(
    T* ptr, 
    cl::sycl::range<dimensions>(int range_size)


Create buffers for the 2 input vectors as well as for the output one.

Location in the source code:

// <<Setup device memory>>

// buffer for input 1
sycl::buffer<sycl::float4, 1> buf_a(&a, cl::sycl::range<1>{1});

Access Device Memory - Accessors


SYCL has the concept of accessors which are used to access request control over the device memory from the buffer objects.

The different access modes are encapsulated in the following enum class:

enum class mode { 

that lives in the access namespace.

At this stage we are only going to focus on the more common ones: read, write, and read_write.

A SYCL accessor can be created from a buffer by using the following construct:

// e.g., read access to buffer memory
auto sycl_acc = sycl_buf.get_access<cl::sycl::access::mode::read>(cgh)

where cgh is an instance of the SYCL command group handler object cl::sycl::handler. Accessors are required to be created inside the command group for which they will be used.


Request access from the buffers for access to the data on the device.

  • Read access from the input buffers
  • Write access from the output buffer

Location in the source code:

// <<Request device memory access>>

// read accessor for the memory in `buf_a`
auto a_acc = buf_a.get_access<sycl::access::mode::read>(cgh);

Run it!

Hello World from SYCL
Create your playground on
This playground was created on, our hands-on, knowledge-sharing platform for developers.
Go to
#include <iostream>
// The SYCL header
#include <SYCL/sycl.hpp>
namespace sycl = cl::sycl;
/* This forward declaration will be explained in the next section */
class vector_addition;
int main(int argc, char **) {
// <<Setup host memory>>
// define input vectors
sycl::float4 a = {1.0, 1.0, 1.0, 1.0}; // input 1
sycl::float4 b = {1.0, 1.0, 1.0, 1.0}; // input 2
// define output vector
sycl::float4 c = {0.0, 0.0, 0.0, 0.0}; // output
// <<Setup SYCL queue>>
sycl::queue myQueue(sycl::default_selector{});
// Begin SYCL scope
// <<Setup device memory>>
// define input buffers
// define output buffer
// Submit a command group functor for execution on a queue. This functor
// encapsulates the kernel and the data needed for its execution.
myQueue.submit([&](sycl::handler &cgh) {
// <<Request device memory access>>
// read accessors
// write accessor
cgh.single_task<class vector_addition>([=]() { /* ... */ });
// validation checks
// check buffer via create host accessor
auto host_c = buf_c.get_access<sycl::access::mode::read>();
std::cout << "We have successfully set up and accessed device memory!\n";
// End SYCL scope
return 0;
codingame x discord
Join the CodinGame community on Discord to chat about puzzle contributions, challenges, streams, blog articles - all that good stuff!
Online Participants