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.
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.
#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)); }
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:
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
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.
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.