C and C++ Standard Libraries Support

This extension enables you to use a set of functions from the C and C++ standard libraries in SYCL device code.

Function declarations are taken from the standard headers (for example, from <assert.h> or <cassert>), and the corresponding header must be explicitly included in user code.

Implementation requires a special device library to be linked with a SYCL program. The library must match the C or C++ standard library used to compile the program.

For example, on Linux* with GNU glibc:

clang++ -fsycl -c main.cpp -o main.o
clang++ -fsycl main.o $(SYCL_INSTALL)/lib/libsycl-glibc.o -o a.out

Or on Windows*:

clang++ -fsycl -c main.cpp -o main.obj
clang++ -fsycl main.obj %SYCL_INSTALL%/lib/libsycl-msvc.o -o a.exe

Supported functions from the C standard library:

Function Header
assert macro <assert.h> or <cassert>
logf, log <math.h> or <cmath>
expf, exp <math.h> or <cmath>
frexpf, frexp <math.h> or <cmath>
ldexpf, ldexp <math.h> or <cmath>
log10f, log10 <math.h> or <cmath>
modff, modf <math.h> or <cmath>
exp2f, exp2 <math.h> or <cmath>
expm1f, expm1 <math.h> or <cmath>
ilogbf, ilogb <math.h> or <cmath>
log1pf, log1p <math.h> or <cmath>
log2f, log2 <math.h> or <cmath>
logbf, logb <math.h> or <cmath>
sqrtf, sqrt <math.h> or <cmath>
cbrtf, cbrt <math.h> or <cmath>
hypotf, hypot <math.h> or <cmath>
erff, erf <math.h> or <cmath>
erfcf, erfc <math.h> or <cmath>
tgammaf, tgamma <math.h> or <cmath>
lgammaf, lgamma <math.h> or <cmath>
fmodf, fmod <math.h> or <cmath>
remainderf, remainder <math.h> or <cmath>
remquof, remquo <math.h> or <cmath>
nextafterf, nextafter <math.h> or <cmath>
fdimf, fdim <math.h> or <cmath>
fmaf, fma <math.h> or <cmath>
sinf, sin <math.h> or <cmath>
cosf, cos <math.h> or <cmath>
tanf, tan <math.h> or <cmath>
powf, pow <math.h> or <cmath>
acosf, acos <math.h> or <cmath>
asinf, asin <math.h> or <cmath>
atanf, atan <math.h> or <cmath>
atan2f, atan2 <math.h> or <cmath>
coshf, cosh <math.h> or <cmath>
sinhf, sinh <math.h> or <cmath>
tanhf, tanh <math.h> or <cmath>
acoshf, acosh <math.h> or <cmath>
asinhf, asinh <math.h> or <cmath>
atanhf, atanh <math.h> or <cmath>
cimagf, cimag <complex.h>
crealf, creal <complex.h>
cargf, carg <complex.h>
cabsf, cabs <complex.h>
cprojf, cproj <complex.h>
cexpf, cexp <complex.h>
clogf, clog <complex.h>
cpowf, cpow <complex.h>
cpolarf, cpolar <complex.h>
csqrtf, csqrt <complex.h>
csinhf, csinh <complex.h>
ccoshf, ccosh <complex.h>
ctanhf, ctanh <complex.h>
csinf, csin <complex.h>
ccosf, ccos <complex.h>
ctanf, ctan <complex.h>
casinhf, casinh <complex.h>
cacoshf, cacosh <complex.h>
catanhf, catanh <complex.h>
casinf, casin <complex.h>
cacosf, cacos <complex.h>
catanf, catan <complex.h>

All functions are grouped into different device libraries based on functionalities. C and C++ standard library groups functions and classes by purpose (for example, <math.h> for mathematical operations and transformations), and device library infrastructure uses this as a baseline.

Note

Only the GNU glibc, Microsoft* C libraries are currently supported. The device libraries for <math.h> and <complex.h> are ready for Linux* and Windows* - support will be added in the future. Not all functions from <math.h> are supported right now, the following math functions are not currently supported:

Device libraries can't support both single and double precision because some underlying devices may not support double precision.

Usage Example

#include <assert.h>
#include <CL/sycl.hpp>

template <typename T, size_t N>
void simple_vadd(const std::array<T, N>& VA, const std::array<T, N>& VB,
                 std::array<T, N>& VC) {
  // ...
  cl::sycl::range<1> numOfItems{N};
  cl::sycl::buffer<T, 1> bufferA(VA.data(), numOfItems);
  cl::sycl::buffer<T, 1> bufferB(VB.data(), numOfItems);
  cl::sycl::buffer<T, 1> bufferC(VC.data(), numOfItems);

  deviceQueue.submit([&](cl::sycl::handler& cgh) {
    auto accessorA = bufferA.template get_access<sycl_read>(cgh);
    auto accessorB = bufferB.template get_access<sycl_read>(cgh);
    auto accessorC = bufferC.template get_access<sycl_write>(cgh);

    cgh.parallel_for<class SimpleVadd<T>>(numOfItems,
    [=](cl::sycl::id<1> wiID) {
        accessorC[wiID] = accessorA[wiID] + accessorB[wiID];
        assert(accessorC[wiID] > 0 && "Invalid value");
    });
  });
  deviceQueue.wait_and_throw();
}
#include <math.h>
#include <CL/sycl.hpp>

void device_sin_test() {
  cl::sycl::queue deviceQueue;
  cl::sycl::range<1> numOfItems{1};
  float  result_f = -1.f;
  double result_d = -1.d;
  {
    cl::sycl::buffer<float, 1> buffer1(&result_f, numOfItems);
    cl::sycl::buffer<double, 1> buffer2(&result_d, numOfItems);
    deviceQueue.submit([&](cl::sycl::handler &cgh) {
      auto res_access1 = buffer1.get_access<sycl_write>(cgh);
      auto res_access2 = buffer2.get_access<sycl_write>(cgh);
      cgh.single_task<class DeviceSin>([=]() {
        res_access1[0] = sinf(0.f);
        res_access2[0] = sin(0.0);
      });
    });
  }
  assert((result_f == 0.f) && (result_d == 0.0));
}

Front End

Once the system header is included, the corresponding functions can be used in SYCL device code. This results in a handful of unresolved functions in the LLVM IR after Clang:

; Function Attrs: noreturn nounwind
declare dso_local spir_func void @__assert_fail(
    i8 addrspace(4)*,
    i8 addrspace(4)*,
    i32, i8 addrspace(4)*
)

[...]
cond.false:
  call spir_func void @__assert_fail([...])
  unreachable

The C and C++ specifications do not define names and signatures of the functions from libc implementation that are used for a particular function. For example, the assert macro expands to the following according to the library:

This makes it difficult to handle all possible cases in device compilers. In order to facilitate porting to new platforms, and to avoid imposing a lot of boilerplate code in every device compiler, wrapper libraries are provided with the SYCL compiler that "lower" libc implementation-specific functions into a stable set of functions, which can be later handled by a device compiler.

clang++ -fsycl -c main.cpp -o main.o
clang++ -fsycl main.o $(SYCL_INSTALL)/lib/libsycl-glibc.o -o a.out

This libsycl-glibc.o is one of these wrapper libraries: it provides definitions for a glibc-specific library function, and these definitions call the corresponding functions from the __devicelib_* set of functions.

For example, __assert_fail from IR above gets transformed as follows:

; Function Attrs: noreturn nounwind
declare dso_local spir_func void @__devicelib_assert_fail(
    i8 addrspace(4)*,
    i8 addrspace(4)*,
    i32,
    i8 addrspace(4)*
)

; Function Attrs: noreturn nounwind
define dso_local spir_func void @__assert_fail(
    i8 addrspace(4)*,
    i8 addrspace(4)*,
    i32,
    i8 addrspace(4)*
) {
    call spir_func void @__devicelib_assert_fail([...])
}

[...]
cond.false:
  call spir_func void @__assert_fail([...])
  unreachable

A single wrapper object provides function wrappers for all supported library functions. Every supported C library implementation (MSVC or glibc) has its own wrapper library object:

SPIR-V

Standard library functions are represented as external (import) functions in SPIR-V:

8 Decorate 67 LinkageAttributes "__devicelib_assert_fail" Import
...
2 Label 846
8 FunctionCall 63 864 67 855 857 863 859
1 Unreachable

Device Compiler

The device compiler is free to implement these __devicelib_* functions. In order to indicate support for a particular set of functions, the underlying runtime must support the corresponding OpenCL (PI) extension.

Fallback Implementation

If a device compiler does not indicate "native" support for a particular function, a fallback library is linked "just in time" (JIT) by the SYCL Runtime. This library is distributed with the SYCL Runtime and resides in the same directory as libsycl.so or sycl.dll.

A fallback library is implemented as a device-agnostic SPIR-V program, and it is intended to work for any device that supports SPIR-V.

Every set of functions is implemented in a separate fallback library. For example, a fallback for the cl_intel_devicelib_cassert extension is provided as libsycl-fallback-cassert.spv.

Note that "ahead-of-time" (AOT) compilation is not yet supported. The driver must check for extension support and link the corresponding SPIR-V fallback implementation, but this is not yet implemented.