The SYCL 1.2 specification defines a buffered output stream class, that is based on the std::ostream of C++. This class aims to replace the C style printf of OpenCL C, which enables outputting information via stdout from within a SYCL kernel. The stream class provides a series of stream operators for basic and SYCL types as well as a series of manipulators to make logging from kernels easier. As the stream class is modelled off of the std::ostream of C++ it supports many common manipulators that are the equivalent of that of C++ and follows the same naming conventions. The stream class is designed for use only within kernels in order to stream from the kernel back to the host CPU for printing via stdout and is passed to a kernel in a similar fashion to accessors.

The stream object (cl::sycl::stream) is constructed on the host within the command group scope in the same manner as an accessor with a handler instance. When constructing a stream object, the buffer size and the maximum statement length must be specified. The buffer size is the size in bytes of the character buffer that is used internally to store the stream output. The maximum statement length is the maximum number of characters that can be used in a single statement on the stream. Note that formatting of non-basic types would add extra characters to provide a human-readable output, hence, the total number of characters per statement or the overall size may be larger than expected.

/* Create a stream object with a buffer of 1024 and a maximum statement length of 128. */
cl::sycl::stream os(1024, 128, handler);

For example, the following statement on the kernel would require a buffer size of 58 characters, and a statement 32. Note that the end of line requires two characters.

int number = 33;
os << "This is a line with a number: " << number;
os << "This is a different line" << cl::sycl::endl;

The stream class has stream operators for all the basic types. In addition, the stream class also provides additional stream operators for all variations of the vec class. Index-space classes are still to be implemented (e.g item, nd_item, …​).

The stream class provides some of the manipulators that are available for the std::ostream: endl, scientific and hex. The following are still to be implemented: oct, setprecision, setbase, showbase.

The listing below shows several possible use cases of the stream class with different operators.

/* check the stream operator for bool. */
os << true << endl;

/* check the stream operator for char. */
os << char('A') << endl;

/* check the stream operator for const char *. */
os << "hello world" << endl;

/* check the stream operator for positive int. */
os << int(1234) << endl;

/* check the stream operator for negative int. */
os << int(-1234) << endl;

/* check the stream operator for unsigned int. */
os << unsigned(1234) << endl;

/* check the stream operator for positive float. */
os << 1234.5678f << endl;

/* check the stream operator for negative float. */
os << -1234.5678f << endl;

/* check the stream operator for a pointer. */
int val = 0;
os << &val << endl;

/* check the stream operator for a int2. */
os << int2(12, -7) << endl;

/* check the stream operator for a float4. */
os << float4(0.43f, -1.27f, 0.111f, -0.89f) << endl;

/* check the stream operator for a char16. */
os << char16('H', 'e', 'l', 'l', 'o', ' ', 'W', 'o', 'r', 'l', 'd', ' ', '!', ' ', ':', ')') << endl;

/* check scientific stream manipulator. */
os << scientific << 1234.5678f << endl;

/* check hex stream manipulator. */
os << hex << 1234 << endl;

When strings are stream to the object within a kernel, the contents are stored on an internal cl_mem object until the kernel finishes execution. The output is then output to stdout when the stream object is destroyed by the runtime. Note that there is no guarantee on when the output will be displayed on stdout, only that it will be printed when the command group is executed.

• The current implementation is still a work in progress, the majority of the feature set is provided however some of the stream operators and manipulators from the specification are not yet implemented.

• Operations to streams are not atomic, note that the output of various streams may get interleaved.