[SYCL] Implement device libraries: C assert (#922)

Device libraries provide a way to support functions from standard C
and C++ system headers in SYCL device code.

Runtime library requirements (functions that are normally defined in
libc, libm or libstdc++) are provided in a SPIR-V
library (libsycl-fallback.spv), which is linked at JIT (or AOT) time.

If a particular device supports a library OpenCL
extension (e.g. cl_intel_devicelib_assert), then SPIR-V implementation
is not linked and the device compiler has to provide definitions for
the corresponding functions.

Signed-off-by: Andrew Savonichev <[email protected]>

Author: Andrew Savonichev

sycl/doc/SYCLEnvironmentVariables.md

@@ -18,6 +18,7 @@ subject to change. Do not rely on these variables in production code.
 | SYCL_DUMP_IMAGES | Any(*) | Dump device image binaries to file. Control has no effect if SYCL_USE_KERNEL_SPV is set. |
 | SYCL_PRINT_EXECUTION_GRAPH | Described [below](#sycl_print_execution_graph-options) | Print execution graph to DOT text file. |
 | SYCL_THROW_ON_BLOCK | Any(*) | Throw an exception on attempt to wait for a blocked command.  |
+| SYCL_DEVICELIB_INHIBIT_NATIVE | String of device library extensions (separated by a whitespace) | Do not rely on device native support for devicelib extensions listed in this option. |
 
 `(*) Note: Any means this environment variable is effective when set to any non-null value.`
 

sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst

@@ -0,0 +1,162 @@
+C and C++ Standard libraries support
+===================================
+
+This extension enables a set of functions from the C and C++ standard
+libraries, and allows to use them in SYCL device code.
+
+Function declarations are taken from the standard headers (e.g. from
+<assert.h> or <cassert>), and the corresponding header has to be
+explicitly included in user code.
+
+Implementation requires a special device library to be linked with a
+SYCL program. The library should match the C or C++ standard library
+used to compile the program:
+
+For example, on Linux with GNU glibc:
+.. code:
+   clang++ -fsycl -c main.cpp -o main.o
+   clang++ -fsycl main.o $(SYCL_INSTALL)/lib/libsycl-glibc.o -o a.out
+
+or, in case of Windows:
+.. code:
+   clang++ -fsycl -c main.cpp -o main.obj
+   clang++ -fsycl main.obj %SYCL_INSTALL%/lib/libsycl-msvc.o -o a.exe
+
+List of supported functions from C standard library:
+  - assert macro (from <assert.h> or <cassert>)
+
+NOTE: only the GNU glibc and Microsoft C libraries are currently
+supported.
+
+Example of usage
+================
+
+.. code: c++
+   #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();
+   }
+
+Frontend
+========
+
+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 LLVM IR after clang:
+
+.. code:
+    ; 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:
+
+  - in Glibc and musl libraries it expands to `__assert_fail`
+  - in MSVC library it expands to `_wassert`
+  - in newlib library it expands to `__assert_func`
+
+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,
+that can be later handled by a device compiler.
+
+.. code:
+   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 glibc specific library function, and these definitions
+call the corresponding functions from `__devicelib_*` set of
+functions.
+
+For example, `__assert_fail` from IR above gets transformed into:
+.. code:
+    ; 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:
+
+  - libsycl-glibc.o
+  - libsycl-msvc.o
+
+SPIR-V
+======
+
+Standard library functions are represented as external (import)
+functions in SPIR-V:
+
+.. code:
+   8 Decorate 67 LinkageAttributes "__devicelib_assert_fail" Import
+   ...
+   2 Label 846
+   8 FunctionCall 63 864 67 855 857 863 859
+   1 Unreachable
+
+Device compiler
+===============
+
+Device compiler is free to implement these `__devicelib_*` functions.
+In order to indicate support for a particular set of functions,
+underlying runtime have to support the corresponding OpenCL (PI)
+extension. See ``DeviceLibExtensions.rst`` for a list of supported
+functions and corresponding extensions.
+
+Fallback implementation
+=======================
+
+If a device compiler does not indicate "native" support for a
+particular function, a fallback library is linked at JIT time by the
+SYCL Runtime. This library is distributed with the SYCL Runtime and
+resides in the same directory as the `libsycl.so` or `sycl.dll`.
+
+A fallback library is implemented as a device-agnostic SPIR-V program,
+and it is supposed 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 `cl_intel_devicelib_cassert`
+extension is provided as `libsycl-fallback-cassert.spv`
+
+NOTE that AOT compilation is not yet supported. Driver will have to
+check for extension support and link the corresponding SPIR-V fallback
+implementation, but this is not implemented yet.

sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst

@@ -0,0 +1,34 @@
+Device library extensions
+===================================
+
+Device compiler that indicates support for a particular extension is
+supposed to support *all* the corresponding functions.
+
+cl_intel_devicelib_cassert
+==========================
+
+.. code:
+   void __devicelib_assert_fail(__generic const char *expr,
+                                __generic const char *file,
+                                int32_t line,
+                                __generic const char *func,
+                                size_t gid0, size_t gid1, size_t gid2,
+                                size_t lid0, size_t lid1, size_t lid2);
+Semantic:
+the function is called when an assertion expression `expr` is false,
+and it indicates that a program does not execute as expected.
+The function should print a message containing the information
+provided in the arguments. In addition to that, the function is free
+to terminate the current kernel invocation.
+
+Arguments:
+
+  - `expr` is a string representation of the assertion condition
+  - `file` and `line` are the source code location of the assertion
+  - `func` (optional, may be NULL)  name of the function containing the assertion
+  - `gidX` current work-item global id
+  - `lidX` current work-item local id
+
+Example of a message:
+.. code:
+   foo.cpp:42: void foo(int): global id: [0,0,0], local id: [0,0,0] Assertion `buf[wiID] == 0 && "Invalid value"` failed.

sycl/include/CL/sycl/detail/context_impl.hpp

@@ -12,6 +12,7 @@
 #include <CL/sycl/detail/os_util.hpp>
 #include <CL/sycl/detail/pi.hpp>
 #include <CL/sycl/detail/platform_impl.hpp>
+#include <CL/sycl/detail/program_manager/program_manager.hpp>
 #include <CL/sycl/detail/usm_dispatch.hpp>
 #include <CL/sycl/exception_list.hpp>
 #include <CL/sycl/info/info_desc.hpp>
@@ -103,6 +104,12 @@ class context_impl {
   /// @return an instance of raw plug-in context handle.
   const RT::PiContext &getHandleRef() const;
 
+  /// Unlike `get_info<info::context::devices>', this function returns a
+  /// reference.
+  const vector_class<device> &getDevices() const {
+    return MDevices;
+  }
+
   /// Gets cached programs.
   ///
   /// @return a map of cached programs.
@@ -122,6 +129,23 @@ class context_impl {
   ///
   /// @return a pointer to USM dispatcher.
   std::shared_ptr<usm::USMDispatcher> getUSMDispatch() const;
+
+  /// In contrast to user programs, which are compiled from user code, library
+  /// programs come from the SYCL runtime. They are identified by the
+  /// corresponding extension:
+  ///
+  ///  cl_intel_devicelib_assert -> #<pi_program with assert functions>
+  ///  cl_intel_devicelib_complex -> #<pi_program with complex functions>
+  ///  etc.
+  ///
+  /// See `doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst' for
+  /// more details.
+  ///
+  /// @returns a map with device library programs.
+  std::map<DeviceLibExt, RT::PiProgram> &getCachedLibPrograms() {
+    return MCachedLibPrograms;
+  }
+
 private:
   async_handler MAsyncHandler;
   vector_class<device> MDevices;
@@ -132,6 +156,7 @@ class context_impl {
   std::map<KernelSetId, RT::PiProgram> MCachedPrograms;
   std::map<RT::PiProgram, std::map<string_class, RT::PiKernel>> MCachedKernels;
   std::shared_ptr<usm::USMDispatcher> MUSMDispatch;
+  std::map<DeviceLibExt, RT::PiProgram> MCachedLibPrograms;
 };
 
 } // namespace detail

sycl/include/CL/sycl/detail/pi.h

@@ -91,7 +91,8 @@ typedef enum {
   PI_DEVICE_INFO_PARTITION_TYPE = CL_DEVICE_PARTITION_TYPE,
   PI_DEVICE_INFO_NAME           = CL_DEVICE_NAME,
   PI_DEVICE_VERSION             = CL_DEVICE_VERSION,
-  PI_DEVICE_MAX_WORK_GROUP_SIZE = CL_DEVICE_MAX_WORK_GROUP_SIZE
+  PI_DEVICE_MAX_WORK_GROUP_SIZE = CL_DEVICE_MAX_WORK_GROUP_SIZE,
+  PI_DEVICE_INFO_EXTENSIONS     = CL_DEVICE_EXTENSIONS
 } _pi_device_info;
 
 // TODO: populate

sycl/include/CL/sycl/detail/program_manager/program_manager.hpp

@@ -41,6 +41,10 @@ using DeviceImage = pi_device_binary_struct;
 // be attempted to de-allocate.
 struct ImageDeleter;
 
+enum DeviceLibExt {
+  cl_intel_devicelib_assert = 0
+};
+
 // Provides single loading and building OpenCL programs with unique contexts
 // that is necessary for no interoperability cases with lambda.
 class ProgramManager {
@@ -70,8 +74,13 @@ class ProgramManager {
 
   DeviceImage &getDeviceImage(OSModuleHandle M, KernelSetId KSId,
                               const context &Context);
-  void build(RT::PiProgram Program, const string_class &Options,
-             std::vector<RT::PiDevice> Devices);
+  using ProgramPtr = unique_ptr_class<remove_pointer_t<RT::PiProgram>,
+                                      decltype(&::piProgramRelease)>;
+  ProgramPtr build(ProgramPtr Program, RT::PiContext Context,
+                   const string_class &Options,
+                   const std::vector<RT::PiDevice> &Devices,
+                   std::map<DeviceLibExt, RT::PiProgram> &CachedLibPrograms,
+                   bool LinkDeviceLibs = false);
   /// Provides a new kernel set id for grouping kernel names together
   KernelSetId getNextKernelSetId() const;
   /// Returns the kernel set associated with the kernel, handles some special

sycl/include/CL/sycl/device.hpp

@@ -157,6 +157,12 @@ class device {
 
   template <class Obj>
   friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject);
+
+  template <class T>
+  friend
+      typename std::add_pointer<typename decltype(T::impl)::element_type>::type
+      detail::getRawSyclObjImpl(const T &SyclObject);
+
   template <class T>
   friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj);
 };

sycl/source/CMakeLists.txt

@@ -143,3 +143,4 @@ install(TARGETS ${SYCL_RT_LIBS}
     ARCHIVE DESTINATION "lib" COMPONENT sycl
     LIBRARY DESTINATION "lib" COMPONENT sycl
     RUNTIME DESTINATION "bin" COMPONENT sycl)
+add_subdirectory(detail/devicelib)

sycl/source/detail/context_impl.cpp

@@ -90,6 +90,10 @@ context_impl::~context_impl() {
       PI_CALL(piKernelRelease)(KernIt.second);
     PI_CALL(piProgramRelease)(ToBeDeleted);
   }
+  for (auto LibProg : MCachedLibPrograms) {
+    assert(LibProg.second && "Null program must not be kept in the cache");
+    PI_CALL(piProgramRelease)(LibProg.second);
+  }
 }
 
 const async_handler &context_impl::get_async_handler() const {

sycl/source/detail/devicelib/CMakeLists.txt

@@ -0,0 +1,66 @@
+# Place device libraries near the libsycl.so library in a build
+# directory
+if (WIN32)
+    set(binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}")
+else()
+    set(binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}")
+endif()
+
+set(clang $<TARGET_FILE:clang>)
+
+set(compile_opts
+    # suppress an error about SYCL_EXTERNAL
+    -Wno-sycl-strict
+    # for CL/__spirv/spirv_vars.hpp
+    -I${sycl_inc_dir})
+
+if (WIN32)
+    set(devicelib-obj-file ${binary_dir}/libsycl-msvc.o)
+    add_custom_command(OUTPUT ${devicelib-obj-file}
+                       COMMAND ${clang} -fsycl -c
+                               ${compile_opts}
+                               ${CMAKE_CURRENT_SOURCE_DIR}/msvc_wrapper.cpp
+                               -o ${devicelib-obj-file}
+                       MAIN_DEPENDENCY msvc_wrapper.cpp
+                       DEPENDS wrapper.h clang
+                       VERBATIM)
+else()
+    set(devicelib-obj-file ${binary_dir}/libsycl-glibc.o)
+    add_custom_command(OUTPUT ${devicelib-obj-file}
+                       COMMAND ${clang} -fsycl -c
+                               ${compile_opts}
+                               ${CMAKE_CURRENT_SOURCE_DIR}/glibc_wrapper.cpp
+                               -o ${devicelib-obj-file}
+                       MAIN_DEPENDENCY glibc_wrapper.cpp
+                       DEPENDS wrapper.h clang
+                       VERBATIM)
+endif()
+
+add_custom_command(OUTPUT ${binary_dir}/libsycl-fallback-cassert.spv
+                   COMMAND ${clang} -S -fsycl-device-only -fno-sycl-use-bitcode
+                           ${compile_opts}
+                           ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp
+                           -o ${binary_dir}/libsycl-fallback-cassert.spv
+                   MAIN_DEPENDENCY fallback-cassert.cpp
+                   DEPENDS wrapper.h clang llvm-spirv
+                   VERBATIM)
+
+add_custom_target(devicelib-obj DEPENDS ${devicelib-obj-file})
+add_custom_target(devicelib-spv DEPENDS ${binary_dir}/libsycl-fallback-cassert.spv)
+add_dependencies(sycl devicelib-obj devicelib-spv)
+if (MSVC)
+    add_dependencies(sycld devicelib-obj devicelib-spv)
+endif()
+
+# Place device libraries near the libsycl.so library in an install
+# directory as well
+if (WIN32)
+    set(install_dest bin)
+else()
+    set(install_dest lib)
+endif()
+
+install(FILES ${devicelib-obj-file}
+              ${binary_dir}/libsycl-fallback-cassert.spv
+        DESTINATION ${install_dest}
+        COMPONENT sycl)