From 9c588433c4dcace2eca19385dba570124c3d474f Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 29 Aug 2019 12:43:00 +0300 Subject: [PATCH 01/32] [SYCL] Implement device libraries: C assert 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 --- .../C-CXX-StandardLibrary.rst | 160 ++++++++++++++ .../DeviceLibExtensions.rst | 34 +++ sycl/include/CL/sycl/detail/context_impl.hpp | 16 ++ sycl/include/CL/sycl/detail/pi.h | 3 +- .../program_manager/program_manager.hpp | 8 +- sycl/source/CMakeLists.txt | 1 + sycl/source/detail/devicelib/CMakeLists.txt | 65 ++++++ .../detail/devicelib/fallback-cassert.cpp | 38 ++++ .../source/detail/devicelib/glibc_wrapper.cpp | 28 +++ sycl/source/detail/devicelib/msvc_wrapper.cpp | 46 ++++ sycl/source/detail/devicelib/wrapper.h | 21 ++ .../program_manager/program_manager.cpp | 174 ++++++++++++++- sycl/test/devicelib/assert-windows.cpp | 69 ++++++ sycl/test/devicelib/assert.cpp | 203 ++++++++++++++++++ sycl/test/lit.cfg.py | 1 + 15 files changed, 855 insertions(+), 12 deletions(-) create mode 100644 sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst create mode 100644 sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst create mode 100644 sycl/source/detail/devicelib/CMakeLists.txt create mode 100644 sycl/source/detail/devicelib/fallback-cassert.cpp create mode 100644 sycl/source/detail/devicelib/glibc_wrapper.cpp create mode 100644 sycl/source/detail/devicelib/msvc_wrapper.cpp create mode 100644 sycl/source/detail/devicelib/wrapper.h create mode 100644 sycl/test/devicelib/assert-windows.cpp create mode 100644 sycl/test/devicelib/assert.cpp diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst new file mode 100644 index 0000000000000..30a2d156590e8 --- /dev/null +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst @@ -0,0 +1,160 @@ +C and C++ Standard libraries support +=================================== + +This extension enables a set of functions from 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 +or ), and the corresponding header has to be explicitly +included in user code. + +List of supported functions from C standard library: + - assert macro (from assert.h) + +NOTE: only the GNU glibc and Microsoft C libraries are currently +supported. + +Device library is distributed with the compiler, and it has to be +explicitly linked by a user. + +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 + +Example of usage +================ + +.. code: c++ + #include + #include + + template + void simple_vadd(const std::array& VA, const std::array& VB, + std::array& VC) { + // ... + cl::sycl::range<1> numOfItems{N}; + cl::sycl::buffer bufferA(VA.data(), numOfItems); + cl::sycl::buffer bufferB(VB.data(), numOfItems); + cl::sycl::buffer bufferC(VC.data(), numOfItems); + + deviceQueue.submit([&](cl::sycl::handler& cgh) { + auto accessorA = bufferA.template get_access(cgh); + auto accessorB = bufferB.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + cgh.parallel_for>(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 + +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, a +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) have 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. diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst new file mode 100644 index 0000000000000..6596dd5b16ac9 --- /dev/null +++ b/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 +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 a 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): local id: [0,0,0], global id: [0,0,0] Assertion `buf[wiID] == 0 && "Invalid value"` failed. diff --git a/sycl/include/CL/sycl/detail/context_impl.hpp b/sycl/include/CL/sycl/detail/context_impl.hpp index 2c76ffccff5ca..be4803a30388f 100644 --- a/sycl/include/CL/sycl/detail/context_impl.hpp +++ b/sycl/include/CL/sycl/detail/context_impl.hpp @@ -121,6 +121,21 @@ class context_impl { /// /// @return a pointer to USM dispatcher. std::shared_ptr getUSMDispatch() const; + + /// Returns device library programs: in contrast to user programs, which are a + /// part of user code, these programs come from the SYCL runtime. Device + /// libraries are identified by the corresponding extension name: + /// + /// "cl_intel_devicelib_assert" -> # + /// "cl_intel_devicelib_complex" -> # + //// etc. + /// + /// See `doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst' for + /// more details. + std::map &getCachedLibPrograms() { + return MCachedLibPrograms; + } + private: async_handler MAsyncHandler; vector_class MDevices; @@ -131,6 +146,7 @@ class context_impl { std::map MCachedPrograms; std::map> MCachedKernels; std::shared_ptr MUSMDispatch; + std::map MCachedLibPrograms; }; } // namespace detail diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 93e55e5f8b61b..e2974430d7a7e 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/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 diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index 3ebf0fda48088..67da239e3a4b4 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp @@ -70,8 +70,12 @@ class ProgramManager { DeviceImage &getDeviceImage(OSModuleHandle M, KernelSetId KSId, const context &Context); - void build(RT::PiProgram Program, const string_class &Options, - std::vector Devices); + RT::PiProgram build(RT::PiProgram Program, + RT::PiContext Context, + const string_class &Options, + std::vector Devices, + std::map &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 diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 7ecd155a3849b..819162e873ad0 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -137,3 +137,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) diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt new file mode 100644 index 0000000000000..92966b2566698 --- /dev/null +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -0,0 +1,65 @@ +# Place device libraries near the libsycl.so library in a build +# directory +if (MSVC) + set(binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") +else() + set(binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") +endif() + +set(clang $) + +if (MSVC) + set(devicelib-obj-file ${binary_dir}/libsycl-msvc.o) + add_custom_command(OUTPUT ${devicelib-obj-file} + COMMAND ${clang} -fsycl -c + # suppress an error about SYCL_EXTERNAL + -Wno-error=sycl-strict -Wno-sycl-strict + ${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 + # suppress an error about SYCL_EXTERNAL + -Wno-error=sycl-strict -Wno-sycl-strict + ${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} -fsycl-device-only -S -Xclang -emit-llvm-bc + # suppress an error about SYCL_EXTERNAL + -Wno-error=sycl-strict -Wno-sycl-strict + ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp + -o libsycl-fallback-cassert.bc + COMMAND llvm-spirv libsycl-fallback-cassert.bc + -o ${binary_dir}/libsycl-fallback-cassert.spv + MAIN_DEPENDENCY fallback-cassert.cpp + DEPENDS wrapper.h clang + 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 (MSVC) + 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) diff --git a/sycl/source/detail/devicelib/fallback-cassert.cpp b/sycl/source/detail/devicelib/fallback-cassert.cpp new file mode 100644 index 0000000000000..619b3921722db --- /dev/null +++ b/sycl/source/detail/devicelib/fallback-cassert.cpp @@ -0,0 +1,38 @@ +//==--- fallback-cassert.cpp - device agnostic implementation of C assert --==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include "wrapper.h" + +// __attribute((format(...))) enables compiler checks for a format string. +int __spirv_ocl_printf(const __attribute__((opencl_constant)) char* fmt, ...) __attribute__((format(printf, 1, 2))); + +static const __attribute__((opencl_constant)) char assert_fmt[] = + "%s:%d: %s: local id: [%lu,%lu,%lu], global id: [%lu,%lu,%lu] " + "Assertion `%s` failed.\n"; + +SYCL_EXTERNAL +extern "C" void __devicelib_assert_fail(const char *expr, const char *file, + int32_t line, const char *func, + size_t gid0, size_t gid1, size_t gid2, + size_t lid0, size_t lid1, size_t lid2) { + // intX_t types are used instead of `int' and `long' because the format string + // is defined in terms of *device* types (OpenCL types): %d matches a 32 bit + // integer, %lu matches a 64 bit unsigned integer. Host `int' and + // `long' types may be different, so we cannot use them. + __spirv_ocl_printf( + assert_fmt, + file, (int32_t)line, + (func) ? func : "", + (uint64_t)gid0, (uint64_t)gid1, (uint64_t)gid2, + (uint64_t)lid0, (uint64_t)lid1, (uint64_t)lid2, + expr); + + // FIXME: call SPIR-V unreachable instead + // volatile int *die = (int *)0x0; + // *die = 0xdead; +} diff --git a/sycl/source/detail/devicelib/glibc_wrapper.cpp b/sycl/source/detail/devicelib/glibc_wrapper.cpp new file mode 100644 index 0000000000000..90bfd36c8defb --- /dev/null +++ b/sycl/source/detail/devicelib/glibc_wrapper.cpp @@ -0,0 +1,28 @@ +//==--- glibc_wrapper.cpp - wrappers for Glibc internal functions ----------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifdef __SYCL_DEVICE_ONLY__ +#include "wrapper.h" + +#include // for __spirv_BuiltInGlobalInvocationId, + // __spirv_BuiltInLocalInvocationId + +extern "C" { + SYCL_EXTERNAL + void __assert_fail(const char *expr, const char *file, + unsigned int line, const char *func) { + __devicelib_assert_fail(expr, file, line, func, + __spirv_BuiltInGlobalInvocationId.x, + __spirv_BuiltInGlobalInvocationId.y, + __spirv_BuiltInGlobalInvocationId.z, + __spirv_BuiltInLocalInvocationId.x, + __spirv_BuiltInLocalInvocationId.y, + __spirv_BuiltInLocalInvocationId.z); + } +} +#endif // __SYCL_DEVICE_ONLY__ diff --git a/sycl/source/detail/devicelib/msvc_wrapper.cpp b/sycl/source/detail/devicelib/msvc_wrapper.cpp new file mode 100644 index 0000000000000..45528c21c9028 --- /dev/null +++ b/sycl/source/detail/devicelib/msvc_wrapper.cpp @@ -0,0 +1,46 @@ +//==--- msvc_wrapper.cpp - wrappers for Microsoft C library functions ------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifdef __SYCL_DEVICE_ONLY__ +#include "wrapper.h" + +#include // for __spirv_BuiltInGlobalInvocationId, + // __spirv_BuiltInLocalInvocationId + +// Truncates a wide (16 or 32 bit) string (wstr) into an ASCII string (str). +// Any non-ASCII characters are replaced by question mark '?'. +static void __truncate_wchar_char_str(const wchar_t *wstr, char* str, size_t str_size) { + str_size -= 1; // reserve for NULL terminator + while (str_size > 0 && *wstr != L'\0') { + wchar_t w = *wstr++; + *str++ = (w > 0 && w < 127) ? (char)w : '?'; + str_size--; + } + *str = '\0'; +} + +extern "C" { + SYCL_EXTERNAL + void _wassert(const wchar_t *wexpr, const wchar_t *wfile, unsigned line) { + // Paths and expressions that are longer than 256 characters are going to be + // truncated. + char file[256]; + __truncate_wchar_char_str(wfile, file, sizeof(file)); + char expr[256]; + __truncate_wchar_char_str(wexpr, expr, sizeof(expr)); + + __devicelib_assert_fail(expr, file, line, /*func=*/nullptr, + __spirv_BuiltInGlobalInvocationId.x, + __spirv_BuiltInGlobalInvocationId.y, + __spirv_BuiltInGlobalInvocationId.z, + __spirv_BuiltInLocalInvocationId.x, + __spirv_BuiltInLocalInvocationId.y, + __spirv_BuiltInLocalInvocationId.z); + } +} +#endif // __SYCL_DEVICE_ONLY__ diff --git a/sycl/source/detail/devicelib/wrapper.h b/sycl/source/detail/devicelib/wrapper.h new file mode 100644 index 0000000000000..6c9ede694e62f --- /dev/null +++ b/sycl/source/detail/devicelib/wrapper.h @@ -0,0 +1,21 @@ +//==--- wrapper.h - declarations for devicelib functions -----*- C++ -*-----==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#ifndef __SYCL_WRAPPERS_H__ +#define __SYCL_WRAPPERS_H__ + +#include +#include + +SYCL_EXTERNAL +extern "C" void __devicelib_assert_fail(const char *expr, const char *file, + int32_t line, const char *func, + size_t gid0, size_t gid1, size_t gid2, + size_t lid0, size_t lid1, size_t lid2); + +#endif // __SYCL_WRAPPERS_H__ diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 81db4ce875b7d..9f0442af9b27f 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -21,6 +21,7 @@ #include #include #include +#include #include #include #include @@ -87,6 +88,18 @@ static RT::PiProgram createSpirvProgram(const RT::PiContext Context, return Program; } +static void getContextDevices(const RT::PiContext &Context, + std::vector &Devices) { + size_t NumDevices = 0; + PI_CALL(piContextGetInfo) + (Context, PI_CONTEXT_INFO_NUM_DEVICES, sizeof(NumDevices), + &NumDevices, nullptr); + Devices.resize(NumDevices); + PI_CALL(piContextGetInfo) + (Context, PI_CONTEXT_INFO_DEVICES, + sizeof(RT::PiDevice) * Devices.size(), &Devices[0], nullptr); +} + DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, const string_class &KernelName, const context &Context) { @@ -180,11 +193,30 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, unique_ptr_class ProgramManaged( Prg, RT::PluginInformation.PiFunctionTable.piProgramRelease); - build(ProgramManaged.get(), Img.BuildOptions, {}); - RT::PiProgram Program = ProgramManaged.release(); - CachedPrograms[KSId] = Program; + // Link a fallback implementation of device libraries if they are not + // supported by a device compiler. + // Pre-compiled programs are supposed to be already linked. + bool LinkDeviceLibs = + getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; + + std::vector Devices; + getContextDevices(getRawSyclObjImpl(Context)->getHandleRef(), Devices); + + RT::PiProgram BuiltProgram = + build(ProgramManaged.get(), + getRawSyclObjImpl(Context)->getHandleRef(), + Img.BuildOptions, + Devices, + getRawSyclObjImpl(Context)->getCachedLibPrograms(), + LinkDeviceLibs); + CachedPrograms[KSId] = BuiltProgram; + // FIXME: better to replace w/ unique_ptr + if (BuiltProgram != ProgramManaged.get()) { + ProgramManaged.release(); + } - return Program; + + return BuiltProgram; } RT::PiKernel ProgramManager::getOrCreateKernel(OSModuleHandle M, @@ -240,6 +272,72 @@ string_class ProgramManager::getProgramBuildLog(const RT::PiProgram &Program) { return Log; } +static bool loadDeviceLib(const RT::PiContext &Context, + const char *Name, RT::PiProgram &Prog) { + std::string LibSyclDir = OSUtil::getCurrentDSODir(); + std::ifstream File(LibSyclDir + OSUtil::DirSep + Name, + std::ifstream::in | std::ifstream::binary); + if (!File.good()) { + return false; + } + + File.seekg (0, std::ios::end); + size_t FileSize = File.tellg(); + File.seekg(0, std::ios::beg); + std::vector FileContent(FileSize); + File.read(&FileContent[0], FileSize); + File.close(); + + Prog = createSpirvProgram(Context, + (unsigned char*)&FileContent[0], FileSize); + return Prog != nullptr; +} + +static std::string getDeviceExtensions(const RT::PiDevice &Dev) { + std::string DevExt; + size_t DevExtSize = 0; + PI_CALL(piDeviceGetInfo) + (Dev, PI_DEVICE_INFO_EXTENSIONS, + /*param_value_size=*/ 0, + /*param_value=*/ nullptr, + &DevExtSize); + DevExt.resize(DevExtSize); + PI_CALL(piDeviceGetInfo) + (Dev, PI_DEVICE_INFO_EXTENSIONS, + DevExt.size(), + &DevExt[0], + /*param_value_size_ret=*/ nullptr); + return DevExt; +} + +static RT::PiProgram loadDeviceLibFallback( + const RT::PiContext &Context, + const std::string &Extension, + const char *Opts, + const std::vector &Devices, + std::map &CachedLibPrograms) { + + const char* LibFileName = nullptr; + if (Extension == "cl_intel_devicelib_assert") { + LibFileName = "libsycl-fallback-cassert.spv"; + } else { + throw compile_program_error( + std::string("Unknown device library: ") + Extension); + } + // FIXME: cache with respect to compile options + RT::PiProgram &LibProg = CachedLibPrograms[LibFileName]; + bool ShouldCompile = !LibProg; + if (!LibProg && !loadDeviceLib(Context, LibFileName , LibProg)) + throw compile_program_error(std::string("Failed to load ") + LibFileName); + + if (ShouldCompile) + PI_CALL(piProgramCompile) + (LibProg, Devices.size(), Devices.data(), Opts, + 0, nullptr, nullptr, nullptr, nullptr); + + return LibProg; +} + struct ImageDeleter { void operator()(DeviceImage *I) { delete[] I->BinaryStart; @@ -332,8 +430,13 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, return *Img; } -void ProgramManager::build(RT::PiProgram Program, const string_class &Options, - std::vector Devices) { +RT::PiProgram ProgramManager::build( + RT::PiProgram Program, + RT::PiContext Context, + const string_class &Options, + std::vector Devices, + std::map &CachedLibPrograms, + bool LinkDeviceLibs) { if (DbgProgMgr > 0) { std::cerr << ">>> ProgramManager::build(" << Program << ", " << Options @@ -351,9 +454,62 @@ void ProgramManager::build(RT::PiProgram Program, const string_class &Options, if (!Opts) Opts = Options.c_str(); - if (PI_CALL_NOCHECK(piProgramBuild)(Program, Devices.size(), Devices.data(), - Opts, nullptr, nullptr) == PI_SUCCESS) - return; + + std::vector LinkPrograms; + if (LinkDeviceLibs) { + // TODO: SYCL compiler should generate a list of required extensions for a + // particular program in order to allow us do a more fine-grained check here. + // Require *all* possible devicelib extensions for now. + const char* RequiredDeviceLibExt[] = { + "cl_intel_devicelib_assert" + }; + + std::vector DevExtensions(Devices.size()); + for (size_t i = 0; i < Devices.size(); ++i) { + DevExtensions[i] = getDeviceExtensions(Devices[i]); + } + for (const char* Ext : RequiredDeviceLibExt) { + bool InhibitNativeImpl = false; + if (const char* Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { + InhibitNativeImpl = strstr(Env, Ext) != nullptr; + } + + for (const std::string& DevExtList : DevExtensions) { + bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); + if (!DeviceSupports || InhibitNativeImpl) { + LinkPrograms.push_back(loadDeviceLibFallback(Context, Ext, Opts, + Devices, + CachedLibPrograms)); + break; + } + } + } + } + + if (LinkPrograms.empty()) { + pi_result Error = + PI_CALL_NOCHECK(piProgramBuild)(Program, Devices.size(), Devices.data(), + Opts, nullptr, nullptr); + if (Error == PI_SUCCESS) + return Program; + fprintf(stderr, "Error is %d\n", Error); + } else { + // Include the main program and compile/link everything together + PI_CALL(piProgramCompile) + (Program, Devices.size(), Devices.data(), Opts, + 0, nullptr, nullptr, nullptr, nullptr); + LinkPrograms.push_back(Program); + + RT::PiProgram LinkedProg = nullptr; + pi_result Error = PI_CALL_NOCHECK(piProgramLink) + (Context, Devices.size(), Devices.data(), + Opts, LinkPrograms.size(), &LinkPrograms[0], + nullptr, nullptr, &LinkedProg); + if (Error != PI_SUCCESS) { + throw compile_program_error(getProgramBuildLog(Program)); + } + return LinkedProg; + } throw compile_program_error(getProgramBuildLog(Program)); } diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp new file mode 100644 index 0000000000000..2c780fee8b768 --- /dev/null +++ b/sycl/test/devicelib/assert-windows.cpp @@ -0,0 +1,69 @@ +// REQUIRES: cpu,windows +// RUN: %clangxx -fsycl -c %s -o %t.o +// RUN: %clangxx -fsycl %t.o %llvm_build_libs_dir/../bin/libsycl-msvc.o -o %t.out +// +// MSVC implementation of assert does not call an unreachable built-in, so the +// program doesn't terminate when fallback is used. +// +// FIXME: SPIR-V Unreachable should be called from the fallback +// explicitly. Since the test is going to crash, we'll have to follow a similar +// approach as on Linux - call the test in a subprocess. +// +// RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_LINK_FALLBACK=1 CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU %t.out >%t.stdout.fallback 2>%t.stderr.fallback +// RUN: FileCheck %s --check-prefix=CHECK-MESSAGE --input-file %t.stdout.fallback +// CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: : local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. +// +// RUN: FileCheck %s --input-file %t.stdout.fallback --check-prefix=CHECK-FALLBACK +// CHECK-FALLBACK: ---> piProgramLink + +#include +#include +#include + +using namespace cl::sycl; + +constexpr auto sycl_read = cl::sycl::access::mode::read; +constexpr auto sycl_write = cl::sycl::access::mode::write; + +template +void simple_vadd(const std::array& VA, const std::array& VB, + std::array& VC) { + queue deviceQueue([](cl::sycl::exception_list ExceptionList) { + for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { + try { + std::rethrow_exception(ExceptionPtr); + } catch (cl::sycl::exception &E) { + std::cerr << E.what() << std::endl; + } catch (...) { + std::cerr << "Unknown async exception was caught." << std::endl; + } + } + }); + + cl::sycl::range<1> numOfItems{N}; + cl::sycl::buffer bufferA(VA.data(), numOfItems); + cl::sycl::buffer bufferB(VB.data(), numOfItems); + cl::sycl::buffer bufferC(VC.data(), numOfItems); + + deviceQueue.submit([&](cl::sycl::handler& cgh) { + auto accessorA = bufferA.template get_access(cgh); + auto accessorB = bufferB.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + cgh.parallel_for(numOfItems, + [=](cl::sycl::id<1> wiID) { + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + assert(accessorC[wiID] == 0 && "Invalid value"); + }); + }); + deviceQueue.wait_and_throw(); +} + +int main() { + std::array A = {1, 2, 3}; + std::array B = {1, 2, 3}; + std::array C = {0, 0, 0}; + + simple_vadd(A, B, C); + return EXIT_SUCCESS; +} diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp new file mode 100644 index 0000000000000..2b4c69d9d252f --- /dev/null +++ b/sycl/test/devicelib/assert.cpp @@ -0,0 +1,203 @@ +// REQUIRES: cpu,linux +// RUN: %clangxx -fsycl -c %s -o %t.o +// RUN: %clangxx -fsycl %t.o %llvm_build_libs_dir/libsycl-glibc.o -o %t.out +// (see the other RUN lines below; it is a bit complicated) +// +// assert() call in device code guarantees nothing: on some devices it behaves +// in a usual way and terminate a program. On other devices it can print an +// error message and *continue* execution. Less capable devices can even ignore +// an assert! +// +// This makes testing an assert() a bit difficult task, and we have to rely on +// the implementation details to make sure that both "native" and "fallback" +// implementations work as expected. +// +// This test works only on Intel OpenCL CPU implementation, which is known to +// behave as follows: +// +// Fallback mode (aka the best we can do by following the OpenCL spec): +// 1. Assertion condition is printed to *stdout* by the OpenCL printf(). +// 2. Process (both host and device) is terminated by a SIGSEGV. +// +// Native mode (same behavior as libc assert on CPU): +// 1. Assertion condition is printed to *stderr*. +// 2. Process (both host and device) is terminated by a SIGABRT. +// +// Other devices are "covered" by the assert-dummy.cpp test, which doesn't +// verify anything except a successful compilation for a device. +// +// FIXME: assert-dummy.cpp is not implemented yet, so other devices are not +// covered. +// +// How the test works: +// ------------------- +// +// 1. First we verify that a call sequence in SYCL Runtime is correct: +// +// - in the fallback mode we have to link an additional library that +// provides a generic implementation of assert(). +// +// - in the native mode we don't link anything, and call clBuildProgram for +// a user program alone. +// +// 2. Then we test that there is actually a difference between the two +// modes. Since the CPU device is the only device that supports this +// extension natively, we catch the difference between the fallback and the +// native modes: SIGSEGV should occur in the fallback mode, SIGABRT in the +// native mode. +// +// In order to check the signal we fork() and let the child die. Then we +// verify how it was terminated. EXPECTED_SIGNAL environment variable +// controls the expected result. +// +// 3. We also test that a message is printed to the corresponding fd: stdout +// for the fallback mode and stderr for the native mode. In the fallback +// mode the test process dies right after a call to the OpenCL printf(), so +// the message can still be buffered by stdio. We turn the bufferization +// off explicitly. +// +// SYCL_DEVICELIB_LINK_FALLBACK=1 environment variable is used to force a mode +// in SYCL Runtime, so it doesn't look into a device extensions list and always +// link the fallback library. +// +// NOTE that Intel OpenCL CPU Vectorizer crashes when an `unreachable' +// instruction is found in IR. Workaround it for now using +// CL_CONFIG_USE_VECTORIZER=False environment variable. +// +// We also skip the native test entirely (see SKIP_IF_NO_EXT), since the assert +// extension is a new feature and may not be supported by the runtime used with +// SYCL. +// +// Overall this sounds stable enough. What could possibly go wrong? +// +// RUN: env SYCL_PI_TRACE=1 CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGABRT SKIP_IF_NO_EXT=1 %t.out 2>%t.stderr.native >%t.stdout.native +// RUN: FileCheck %s --input-file %t.stdout.native --check-prefixes=CHECK-NATIVE || FileCheck %s --input-file %t.stderr.native --check-prefix CHECK-NOTSUPPORTED +// RUN: FileCheck %s --input-file %t.stderr.native --check-prefixes=CHECK-MESSAGE || FileCheck %s --input-file %t.stderr.native --check-prefix CHECK-NOTSUPPORTED +// +// RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_INHIBIT_NATIVE=cl_intel_devicelib_assert CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU EXPECTED_SIGNAL=SIGSEGV %t.out 2>%t.stderr.fallback >%t.stdout.fallback +// RUN: FileCheck %s --input-file %t.stdout.fallback --check-prefixes=CHECK-FALLBACK,CHECK-MESSAGE +// +// CHECK-NATIVE: ---> piProgramBuild +// CHECK-FALLBACK: ---> piProgramLink +// +// Skip the test if the CPU RT doesn't support the extension yet: +// CHECK-NOTSUPPORTED: Device has no support for cl_intel_devicelib_assert +// +// Anyway, the same message has to be printed for both the fallback and the +// native modes (fallback prints to stdout, while native prints to stderr; we +// already handled this difference in the RUN lines): +// +// CHECK-MESSAGE: {{.*}}assert.cpp:{{[0-9]+}}: auto simple_vadd(const std::array &, const std::array &, std::array &)::(anonymous class)::operator()(cl::sycl::handler &)::(anonymous class)::operator()(cl::sycl::id<1>) const: local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. +// +// Note that the work-item that hits the assert first may vary, since the order +// of execution is undefined. We catch only the first one (whatever id it is). + +#include +#include +#include + +#include +#include +#include +#include + +using namespace cl::sycl; + +constexpr auto sycl_read = cl::sycl::access::mode::read; +constexpr auto sycl_write = cl::sycl::access::mode::write; + +const int EXIT_SKIP_TEST = 42; + +template +void simple_vadd(const std::array& VA, const std::array& VB, + std::array& VC) { + queue deviceQueue([](cl::sycl::exception_list ExceptionList) { + for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { + try { + std::rethrow_exception(ExceptionPtr); + } catch (cl::sycl::exception &E) { + std::cerr << E.what() << std::endl; + } catch (...) { + std::cerr << "Unknown async exception was caught." << std::endl; + } + } + }); + device dev = deviceQueue.get_device(); + bool unsupported = true; + for (auto &ext : dev.get_info()) { + if (ext == "cl_intel_devicelib_assert") { + unsupported = false; + } + } + if (unsupported && getenv("SKIP_IF_NO_EXT")) { + fprintf(stderr, "Device has no support for cl_intel_devicelib_assert, " + "skipping the test\n"); + exit(EXIT_SKIP_TEST); + } + + cl::sycl::range<1> numOfItems{N}; + cl::sycl::buffer bufferA(VA.data(), numOfItems); + cl::sycl::buffer bufferB(VB.data(), numOfItems); + cl::sycl::buffer bufferC(VC.data(), numOfItems); + + deviceQueue.submit([&](cl::sycl::handler& cgh) { + auto accessorA = bufferA.template get_access(cgh); + auto accessorB = bufferB.template get_access(cgh); + auto accessorC = bufferC.template get_access(cgh); + + cgh.parallel_for(numOfItems, + [=](cl::sycl::id<1> wiID) { + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + assert(accessorC[wiID] == 0 && "Invalid value"); + }); + }); + deviceQueue.wait_and_throw(); +} + +int main() { + int child = fork(); + if (child) { + int status = 0; + waitpid(child, &status, 0); + if (WIFEXITED(status) && WEXITSTATUS(status) == EXIT_SKIP_TEST) { + return 0; + } + if (!WIFSIGNALED(status)) { + fprintf(stderr, "error: process did not terminate by a signal\n"); + return 1; + } + int sig = WTERMSIG(status); + int expected = 0; + if (const char* env = getenv("EXPECTED_SIGNAL")) { + if (0 == strcmp(env, "SIGABRT")) { + expected = SIGABRT; + } else if (0 == strcmp(env, "SIGSEGV")) { + expected = SIGSEGV; + } + if (!expected) { + fprintf(stderr, + "EXPECTED_SIGNAL should be set to either \"SIGABRT\", " + "or \"SIGSEGV\"!\n"); + return 1; + } + } + if (sig != expected) { + fprintf(stderr, "error: expected signal %d, got %d\n", expected, sig); + return 1; + } + return 0; + } + + // Turn the bufferization off to not loose the assert message if it is written + // to stdout. + if (setvbuf(stdout, NULL, _IONBF, 0)) { + perror("failed to turn off bufferization on stdout"); + return 1; + } + + std::array A = {1, 2, 3}; + std::array B = {1, 2, 3}; + std::array C = {0, 0, 0}; + + simple_vadd(A, B, C); +} diff --git a/sycl/test/lit.cfg.py b/sycl/test/lit.cfg.py index 27657bbac25f0..b122c0848648f 100644 --- a/sycl/test/lit.cfg.py +++ b/sycl/test/lit.cfg.py @@ -35,6 +35,7 @@ config.test_exec_root = os.path.join(config.sycl_obj_root, 'test') if platform.system() == "Linux": + config.available_features.add('linux') # Propagate 'LD_LIBRARY_PATH' through the environment. if 'LD_LIBRARY_PATH' in os.environ: config.environment['LD_LIBRARY_PATH'] = os.path.pathsep.join((config.environment['LD_LIBRARY_PATH'], config.llvm_build_libs_dir)) From 6a8b3e7fcf2a70351dffbff2046899e57412a5a4 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 16:51:31 +0300 Subject: [PATCH 02/32] Do not use "user" options for device library compilation Signed-off-by: Andrew Savonichev --- .../detail/program_manager/program_manager.cpp | 14 +++++++++----- 1 file changed, 9 insertions(+), 5 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 9f0442af9b27f..8e0ae6a62c0d1 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -313,7 +313,6 @@ static std::string getDeviceExtensions(const RT::PiDevice &Dev) { static RT::PiProgram loadDeviceLibFallback( const RT::PiContext &Context, const std::string &Extension, - const char *Opts, const std::vector &Devices, std::map &CachedLibPrograms) { @@ -324,7 +323,6 @@ static RT::PiProgram loadDeviceLibFallback( throw compile_program_error( std::string("Unknown device library: ") + Extension); } - // FIXME: cache with respect to compile options RT::PiProgram &LibProg = CachedLibPrograms[LibFileName]; bool ShouldCompile = !LibProg; if (!LibProg && !loadDeviceLib(Context, LibFileName , LibProg)) @@ -332,7 +330,14 @@ static RT::PiProgram loadDeviceLibFallback( if (ShouldCompile) PI_CALL(piProgramCompile) - (LibProg, Devices.size(), Devices.data(), Opts, + (LibProg, + // Assume that Devices contains all devices from Context. + Devices.size(), Devices.data(), + // Do not use compile options for library programs: it is not clear + // if user options (image options) are supposed to be applied to + // library program as well, and what actually happens to a SPIR-V + // program if we apply them. + "", 0, nullptr, nullptr, nullptr, nullptr); return LibProg; @@ -477,8 +482,7 @@ RT::PiProgram ProgramManager::build( for (const std::string& DevExtList : DevExtensions) { bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); if (!DeviceSupports || InhibitNativeImpl) { - LinkPrograms.push_back(loadDeviceLibFallback(Context, Ext, Opts, - Devices, + LinkPrograms.push_back(loadDeviceLibFallback(Context, Ext, Devices, CachedLibPrograms)); break; } From bc335d7a80a40902d1e81f5e70d0475570fd3e88 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 19:23:15 +0300 Subject: [PATCH 03/32] Pass unique_ptr with a program to ProgramManager::build() Since build() may return a different program object, we need to properly release the old one (or not, if we return the same object). Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.hpp | 15 +++++---- .../program_manager/program_manager.cpp | 31 +++++++------------ 2 files changed, 21 insertions(+), 25 deletions(-) diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index 67da239e3a4b4..17079efa17024 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp @@ -70,12 +70,15 @@ class ProgramManager { DeviceImage &getDeviceImage(OSModuleHandle M, KernelSetId KSId, const context &Context); - RT::PiProgram build(RT::PiProgram Program, - RT::PiContext Context, - const string_class &Options, - std::vector Devices, - std::map &CachedLibPrograms, - bool LinkDeviceLibs = false); + using ProgramPtr = + unique_ptr_class, + decltype(&::piProgramRelease)>; + ProgramPtr build(ProgramPtr Program, + RT::PiContext Context, + const string_class &Options, + std::vector Devices, + std::map &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 diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 8e0ae6a62c0d1..1093597dc8244 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -189,8 +189,7 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, const DeviceImage &Img = getDeviceImage(M, KSId, Context); RT::PiProgram Prg = createPIProgram(Img, Context); - using PiProgramT = remove_pointer_t; - unique_ptr_class ProgramManaged( + ProgramPtr ProgramManaged( Prg, RT::PluginInformation.PiFunctionTable.piProgramRelease); // Link a fallback implementation of device libraries if they are not @@ -202,21 +201,15 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, std::vector Devices; getContextDevices(getRawSyclObjImpl(Context)->getHandleRef(), Devices); - RT::PiProgram BuiltProgram = - build(ProgramManaged.get(), - getRawSyclObjImpl(Context)->getHandleRef(), - Img.BuildOptions, - Devices, - getRawSyclObjImpl(Context)->getCachedLibPrograms(), - LinkDeviceLibs); - CachedPrograms[KSId] = BuiltProgram; - // FIXME: better to replace w/ unique_ptr - if (BuiltProgram != ProgramManaged.get()) { - ProgramManaged.release(); - } - - - return BuiltProgram; + ProgramPtr BuiltProgram = + build(std::move(ProgramManaged), + getRawSyclObjImpl(Context)->getHandleRef(), + Img.BuildOptions, + Devices, + getRawSyclObjImpl(Context)->getCachedLibPrograms(), + LinkDeviceLibs); + CachedPrograms[KSId] = BuiltProgram.get(); + return BuiltProgram.release(); } RT::PiKernel ProgramManager::getOrCreateKernel(OSModuleHandle M, @@ -444,7 +437,7 @@ RT::PiProgram ProgramManager::build( bool LinkDeviceLibs) { if (DbgProgMgr > 0) { - std::cerr << ">>> ProgramManager::build(" << Program << ", " << Options + std::cerr << ">>> ProgramManager::build(" << Program.get() << ", " << Options << ", ... " << Devices.size() << ")\n"; } const char *Opts = std::getenv("SYCL_PROGRAM_BUILD_OPTIONS"); @@ -492,7 +485,7 @@ RT::PiProgram ProgramManager::build( if (LinkPrograms.empty()) { pi_result Error = - PI_CALL_NOCHECK(piProgramBuild)(Program, Devices.size(), Devices.data(), + PI_CALL_NOCHECK(piProgramBuild)(Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr); if (Error == PI_SUCCESS) return Program; From b03b3e20822683ca4f3887a8a8ab23e6a1351033 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 19:23:56 +0300 Subject: [PATCH 04/32] Refactoring Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.cpp | 99 ++++++++++--------- 1 file changed, 53 insertions(+), 46 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1093597dc8244..0c4f0cbab4d76 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -428,8 +428,41 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, return *Img; } -RT::PiProgram ProgramManager::build( - RT::PiProgram Program, +void getDeviceLibPrograms( + const RT::PiContext Context, + const std::vector Devices, + std::map &CachedLibPrograms, + std::vector &Programs) { + // TODO: SYCL compiler should generate a list of required extensions for a + // particular program in order to allow us do a more fine-grained check here. + // Require *all* possible devicelib extensions for now. + const char* RequiredDeviceLibExt[] = { + "cl_intel_devicelib_assert" + }; + + std::vector DevExtensions(Devices.size()); + for (size_t i = 0; i < Devices.size(); ++i) { + DevExtensions[i] = getDeviceExtensions(Devices[i]); + } + for (const char* Ext : RequiredDeviceLibExt) { + bool InhibitNativeImpl = false; + if (const char* Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { + InhibitNativeImpl = strstr(Env, Ext) != nullptr; + } + + for (const std::string& DevExtList : DevExtensions) { + bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); + if (!DeviceSupports || InhibitNativeImpl) { + Programs.push_back(loadDeviceLibFallback(Context, Ext, Devices, + CachedLibPrograms)); + break; + } + } + } +} + +ProgramManager::ProgramPtr ProgramManager::build( + ProgramPtr Program, RT::PiContext Context, const string_class &Options, std::vector Devices, @@ -455,32 +488,7 @@ RT::PiProgram ProgramManager::build( std::vector LinkPrograms; if (LinkDeviceLibs) { - // TODO: SYCL compiler should generate a list of required extensions for a - // particular program in order to allow us do a more fine-grained check here. - // Require *all* possible devicelib extensions for now. - const char* RequiredDeviceLibExt[] = { - "cl_intel_devicelib_assert" - }; - - std::vector DevExtensions(Devices.size()); - for (size_t i = 0; i < Devices.size(); ++i) { - DevExtensions[i] = getDeviceExtensions(Devices[i]); - } - for (const char* Ext : RequiredDeviceLibExt) { - bool InhibitNativeImpl = false; - if (const char* Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { - InhibitNativeImpl = strstr(Env, Ext) != nullptr; - } - - for (const std::string& DevExtList : DevExtensions) { - bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); - if (!DeviceSupports || InhibitNativeImpl) { - LinkPrograms.push_back(loadDeviceLibFallback(Context, Ext, Devices, - CachedLibPrograms)); - break; - } - } - } + getDeviceLibPrograms(Context, Devices, CachedLibPrograms, LinkPrograms); } if (LinkPrograms.empty()) { @@ -489,26 +497,25 @@ RT::PiProgram ProgramManager::build( Opts, nullptr, nullptr); if (Error == PI_SUCCESS) return Program; - fprintf(stderr, "Error is %d\n", Error); - } else { - // Include the main program and compile/link everything together - PI_CALL(piProgramCompile) - (Program, Devices.size(), Devices.data(), Opts, - 0, nullptr, nullptr, nullptr, nullptr); - LinkPrograms.push_back(Program); - - RT::PiProgram LinkedProg = nullptr; - pi_result Error = PI_CALL_NOCHECK(piProgramLink) - (Context, Devices.size(), Devices.data(), - Opts, LinkPrograms.size(), &LinkPrograms[0], - nullptr, nullptr, &LinkedProg); - if (Error != PI_SUCCESS) { - throw compile_program_error(getProgramBuildLog(Program)); - } - return LinkedProg; + throw compile_program_error(getProgramBuildLog(Program.get())); } - throw compile_program_error(getProgramBuildLog(Program)); + // Include the main program and compile/link everything together + PI_CALL(piProgramCompile) + (Program.get(), Devices.size(), Devices.data(), Opts, + 0, nullptr, nullptr, nullptr, nullptr); + LinkPrograms.push_back(Program.get()); + + RT::PiProgram LinkedProg = nullptr; + pi_result Error = PI_CALL_NOCHECK(piProgramLink) + (Context, Devices.size(), Devices.data(), + Opts, LinkPrograms.size(), &LinkPrograms[0], + nullptr, nullptr, &LinkedProg); + if (Error != PI_SUCCESS) { + throw compile_program_error(getProgramBuildLog(Program.get())); + } + return ProgramPtr(LinkedProg, + RT::PluginInformation.PiFunctionTable.piProgramRelease); } void ProgramManager::addImages(pi_device_binaries DeviceBinary) { From 08a37e3808689bdc38b70a2451b3cef6a368a190 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 19:48:04 +0300 Subject: [PATCH 05/32] Clang format the whole patch Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.hpp | 8 +- .../detail/devicelib/fallback-cassert.cpp | 3 +- .../source/detail/devicelib/glibc_wrapper.cpp | 22 ++- sycl/source/detail/devicelib/msvc_wrapper.cpp | 35 +++-- sycl/source/detail/devicelib/wrapper.h | 2 +- .../program_manager/program_manager.cpp | 125 ++++++++---------- sycl/test/devicelib/assert-windows.cpp | 39 +++--- sycl/test/devicelib/assert.cpp | 51 ++++--- 8 files changed, 133 insertions(+), 152 deletions(-) diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index 17079efa17024..802562ebf81d8 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp @@ -70,11 +70,9 @@ class ProgramManager { DeviceImage &getDeviceImage(OSModuleHandle M, KernelSetId KSId, const context &Context); - using ProgramPtr = - unique_ptr_class, - decltype(&::piProgramRelease)>; - ProgramPtr build(ProgramPtr Program, - RT::PiContext Context, + using ProgramPtr = unique_ptr_class, + decltype(&::piProgramRelease)>; + ProgramPtr build(ProgramPtr Program, RT::PiContext Context, const string_class &Options, std::vector Devices, std::map &CachedLibPrograms, diff --git a/sycl/source/detail/devicelib/fallback-cassert.cpp b/sycl/source/detail/devicelib/fallback-cassert.cpp index 619b3921722db..204741086fdb2 100644 --- a/sycl/source/detail/devicelib/fallback-cassert.cpp +++ b/sycl/source/detail/devicelib/fallback-cassert.cpp @@ -9,7 +9,8 @@ #include "wrapper.h" // __attribute((format(...))) enables compiler checks for a format string. -int __spirv_ocl_printf(const __attribute__((opencl_constant)) char* fmt, ...) __attribute__((format(printf, 1, 2))); +int __spirv_ocl_printf(const __attribute__((opencl_constant)) char *fmt, ...) + __attribute__((format(printf, 1, 2))); static const __attribute__((opencl_constant)) char assert_fmt[] = "%s:%d: %s: local id: [%lu,%lu,%lu], global id: [%lu,%lu,%lu] " diff --git a/sycl/source/detail/devicelib/glibc_wrapper.cpp b/sycl/source/detail/devicelib/glibc_wrapper.cpp index 90bfd36c8defb..403a90cdda378 100644 --- a/sycl/source/detail/devicelib/glibc_wrapper.cpp +++ b/sycl/source/detail/devicelib/glibc_wrapper.cpp @@ -12,17 +12,15 @@ #include // for __spirv_BuiltInGlobalInvocationId, // __spirv_BuiltInLocalInvocationId -extern "C" { - SYCL_EXTERNAL - void __assert_fail(const char *expr, const char *file, - unsigned int line, const char *func) { - __devicelib_assert_fail(expr, file, line, func, - __spirv_BuiltInGlobalInvocationId.x, - __spirv_BuiltInGlobalInvocationId.y, - __spirv_BuiltInGlobalInvocationId.z, - __spirv_BuiltInLocalInvocationId.x, - __spirv_BuiltInLocalInvocationId.y, - __spirv_BuiltInLocalInvocationId.z); - } +extern "C" SYCL_EXTERNAL +void __assert_fail(const char *expr, const char *file, + unsigned int line, const char *func) { + __devicelib_assert_fail(expr, file, line, func, + __spirv_BuiltInGlobalInvocationId.x, + __spirv_BuiltInGlobalInvocationId.y, + __spirv_BuiltInGlobalInvocationId.z, + __spirv_BuiltInLocalInvocationId.x, + __spirv_BuiltInLocalInvocationId.y, + __spirv_BuiltInLocalInvocationId.z); } #endif // __SYCL_DEVICE_ONLY__ diff --git a/sycl/source/detail/devicelib/msvc_wrapper.cpp b/sycl/source/detail/devicelib/msvc_wrapper.cpp index 45528c21c9028..21b430c3ad81e 100644 --- a/sycl/source/detail/devicelib/msvc_wrapper.cpp +++ b/sycl/source/detail/devicelib/msvc_wrapper.cpp @@ -14,7 +14,8 @@ // Truncates a wide (16 or 32 bit) string (wstr) into an ASCII string (str). // Any non-ASCII characters are replaced by question mark '?'. -static void __truncate_wchar_char_str(const wchar_t *wstr, char* str, size_t str_size) { +static void __truncate_wchar_char_str(const wchar_t *wstr, char *str, + size_t str_size) { str_size -= 1; // reserve for NULL terminator while (str_size > 0 && *wstr != L'\0') { wchar_t w = *wstr++; @@ -24,23 +25,21 @@ static void __truncate_wchar_char_str(const wchar_t *wstr, char* str, size_t str *str = '\0'; } -extern "C" { - SYCL_EXTERNAL - void _wassert(const wchar_t *wexpr, const wchar_t *wfile, unsigned line) { - // Paths and expressions that are longer than 256 characters are going to be - // truncated. - char file[256]; - __truncate_wchar_char_str(wfile, file, sizeof(file)); - char expr[256]; - __truncate_wchar_char_str(wexpr, expr, sizeof(expr)); +extern "C" SYCL_EXTERNAL +void _wassert(const wchar_t *wexpr, const wchar_t *wfile, unsigned line) { + // Paths and expressions that are longer than 256 characters are going to be + // truncated. + char file[256]; + __truncate_wchar_char_str(wfile, file, sizeof(file)); + char expr[256]; + __truncate_wchar_char_str(wexpr, expr, sizeof(expr)); - __devicelib_assert_fail(expr, file, line, /*func=*/nullptr, - __spirv_BuiltInGlobalInvocationId.x, - __spirv_BuiltInGlobalInvocationId.y, - __spirv_BuiltInGlobalInvocationId.z, - __spirv_BuiltInLocalInvocationId.x, - __spirv_BuiltInLocalInvocationId.y, - __spirv_BuiltInLocalInvocationId.z); - } + __devicelib_assert_fail(expr, file, line, /*func=*/nullptr, + __spirv_BuiltInGlobalInvocationId.x, + __spirv_BuiltInGlobalInvocationId.y, + __spirv_BuiltInGlobalInvocationId.z, + __spirv_BuiltInLocalInvocationId.x, + __spirv_BuiltInLocalInvocationId.y, + __spirv_BuiltInLocalInvocationId.z); } #endif // __SYCL_DEVICE_ONLY__ diff --git a/sycl/source/detail/devicelib/wrapper.h b/sycl/source/detail/devicelib/wrapper.h index 6c9ede694e62f..6148f0436498b 100644 --- a/sycl/source/detail/devicelib/wrapper.h +++ b/sycl/source/detail/devicelib/wrapper.h @@ -9,8 +9,8 @@ #ifndef __SYCL_WRAPPERS_H__ #define __SYCL_WRAPPERS_H__ -#include #include +#include SYCL_EXTERNAL extern "C" void __devicelib_assert_fail(const char *expr, const char *file, diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 0c4f0cbab4d76..eade277cf5f57 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -92,12 +92,12 @@ static void getContextDevices(const RT::PiContext &Context, std::vector &Devices) { size_t NumDevices = 0; PI_CALL(piContextGetInfo) - (Context, PI_CONTEXT_INFO_NUM_DEVICES, sizeof(NumDevices), - &NumDevices, nullptr); + (Context, PI_CONTEXT_INFO_NUM_DEVICES, sizeof(NumDevices), &NumDevices, + nullptr); Devices.resize(NumDevices); PI_CALL(piContextGetInfo) - (Context, PI_CONTEXT_INFO_DEVICES, - sizeof(RT::PiDevice) * Devices.size(), &Devices[0], nullptr); + (Context, PI_CONTEXT_INFO_DEVICES, sizeof(RT::PiDevice) * Devices.size(), + &Devices[0], nullptr); } DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, @@ -195,19 +195,15 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, // Link a fallback implementation of device libraries if they are not // supported by a device compiler. // Pre-compiled programs are supposed to be already linked. - bool LinkDeviceLibs = - getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; + bool LinkDeviceLibs = getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; std::vector Devices; getContextDevices(getRawSyclObjImpl(Context)->getHandleRef(), Devices); - ProgramPtr BuiltProgram = - build(std::move(ProgramManaged), - getRawSyclObjImpl(Context)->getHandleRef(), - Img.BuildOptions, - Devices, - getRawSyclObjImpl(Context)->getCachedLibPrograms(), - LinkDeviceLibs); + ProgramPtr BuiltProgram = build( + std::move(ProgramManaged), getRawSyclObjImpl(Context)->getHandleRef(), + Img.BuildOptions, Devices, + getRawSyclObjImpl(Context)->getCachedLibPrograms(), LinkDeviceLibs); CachedPrograms[KSId] = BuiltProgram.get(); return BuiltProgram.release(); } @@ -265,24 +261,24 @@ string_class ProgramManager::getProgramBuildLog(const RT::PiProgram &Program) { return Log; } -static bool loadDeviceLib(const RT::PiContext &Context, - const char *Name, RT::PiProgram &Prog) { +static bool loadDeviceLib(const RT::PiContext &Context, const char *Name, + RT::PiProgram &Prog) { std::string LibSyclDir = OSUtil::getCurrentDSODir(); std::ifstream File(LibSyclDir + OSUtil::DirSep + Name, - std::ifstream::in | std::ifstream::binary); + std::ifstream::in | std::ifstream::binary); if (!File.good()) { return false; } - File.seekg (0, std::ios::end); + File.seekg(0, std::ios::end); size_t FileSize = File.tellg(); File.seekg(0, std::ios::beg); std::vector FileContent(FileSize); File.read(&FileContent[0], FileSize); File.close(); - Prog = createSpirvProgram(Context, - (unsigned char*)&FileContent[0], FileSize); + Prog = + createSpirvProgram(Context, (unsigned char *)&FileContent[0], FileSize); return Prog != nullptr; } @@ -290,48 +286,44 @@ static std::string getDeviceExtensions(const RT::PiDevice &Dev) { std::string DevExt; size_t DevExtSize = 0; PI_CALL(piDeviceGetInfo) - (Dev, PI_DEVICE_INFO_EXTENSIONS, - /*param_value_size=*/ 0, - /*param_value=*/ nullptr, - &DevExtSize); + (Dev, PI_DEVICE_INFO_EXTENSIONS, + /*param_value_size=*/0, + /*param_value=*/nullptr, &DevExtSize); DevExt.resize(DevExtSize); PI_CALL(piDeviceGetInfo) - (Dev, PI_DEVICE_INFO_EXTENSIONS, - DevExt.size(), - &DevExt[0], - /*param_value_size_ret=*/ nullptr); + (Dev, PI_DEVICE_INFO_EXTENSIONS, DevExt.size(), &DevExt[0], + /*param_value_size_ret=*/nullptr); return DevExt; } -static RT::PiProgram loadDeviceLibFallback( - const RT::PiContext &Context, - const std::string &Extension, - const std::vector &Devices, - std::map &CachedLibPrograms) { +static RT::PiProgram +loadDeviceLibFallback(const RT::PiContext &Context, + const std::string &Extension, + const std::vector &Devices, + std::map &CachedLibPrograms) { - const char* LibFileName = nullptr; + const char *LibFileName = nullptr; if (Extension == "cl_intel_devicelib_assert") { LibFileName = "libsycl-fallback-cassert.spv"; } else { - throw compile_program_error( - std::string("Unknown device library: ") + Extension); + throw compile_program_error(std::string("Unknown device library: ") + + Extension); } RT::PiProgram &LibProg = CachedLibPrograms[LibFileName]; bool ShouldCompile = !LibProg; - if (!LibProg && !loadDeviceLib(Context, LibFileName , LibProg)) + if (!LibProg && !loadDeviceLib(Context, LibFileName, LibProg)) throw compile_program_error(std::string("Failed to load ") + LibFileName); if (ShouldCompile) PI_CALL(piProgramCompile) - (LibProg, - // Assume that Devices contains all devices from Context. - Devices.size(), Devices.data(), - // Do not use compile options for library programs: it is not clear - // if user options (image options) are supposed to be applied to - // library program as well, and what actually happens to a SPIR-V - // program if we apply them. - "", - 0, nullptr, nullptr, nullptr, nullptr); + (LibProg, + // Assume that Devices contains all devices from Context. + Devices.size(), Devices.data(), + // Do not use compile options for library programs: it is not clear + // if user options (image options) are supposed to be applied to + // library program as well, and what actually happens to a SPIR-V + // program if we apply them. + "", 0, nullptr, nullptr, nullptr, nullptr); return LibProg; } @@ -444,34 +436,33 @@ void getDeviceLibPrograms( for (size_t i = 0; i < Devices.size(); ++i) { DevExtensions[i] = getDeviceExtensions(Devices[i]); } - for (const char* Ext : RequiredDeviceLibExt) { + for (const char *Ext : RequiredDeviceLibExt) { bool InhibitNativeImpl = false; - if (const char* Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { + if (const char *Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { InhibitNativeImpl = strstr(Env, Ext) != nullptr; } - for (const std::string& DevExtList : DevExtensions) { + for (const std::string &DevExtList : DevExtensions) { bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); if (!DeviceSupports || InhibitNativeImpl) { - Programs.push_back(loadDeviceLibFallback(Context, Ext, Devices, - CachedLibPrograms)); + Programs.push_back( + loadDeviceLibFallback(Context, Ext, Devices, CachedLibPrograms)); break; } } } } -ProgramManager::ProgramPtr ProgramManager::build( - ProgramPtr Program, - RT::PiContext Context, - const string_class &Options, - std::vector Devices, - std::map &CachedLibPrograms, - bool LinkDeviceLibs) { +ProgramManager::ProgramPtr +ProgramManager::build(ProgramPtr Program, RT::PiContext Context, + const string_class &Options, + std::vector Devices, + std::map &CachedLibPrograms, + bool LinkDeviceLibs) { if (DbgProgMgr > 0) { - std::cerr << ">>> ProgramManager::build(" << Program.get() << ", " << Options - << ", ... " << Devices.size() << ")\n"; + std::cerr << ">>> ProgramManager::build(" << Program.get() << ", " + << Options << ", ... " << Devices.size() << ")\n"; } const char *Opts = std::getenv("SYCL_PROGRAM_BUILD_OPTIONS"); @@ -492,9 +483,8 @@ ProgramManager::ProgramPtr ProgramManager::build( } if (LinkPrograms.empty()) { - pi_result Error = - PI_CALL_NOCHECK(piProgramBuild)(Program.get(), Devices.size(), Devices.data(), - Opts, nullptr, nullptr); + pi_result Error = PI_CALL_NOCHECK(piProgramBuild)( + Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr); if (Error == PI_SUCCESS) return Program; throw compile_program_error(getProgramBuildLog(Program.get())); @@ -502,15 +492,14 @@ ProgramManager::ProgramPtr ProgramManager::build( // Include the main program and compile/link everything together PI_CALL(piProgramCompile) - (Program.get(), Devices.size(), Devices.data(), Opts, - 0, nullptr, nullptr, nullptr, nullptr); + (Program.get(), Devices.size(), Devices.data(), Opts, 0, nullptr, nullptr, + nullptr, nullptr); LinkPrograms.push_back(Program.get()); RT::PiProgram LinkedProg = nullptr; - pi_result Error = PI_CALL_NOCHECK(piProgramLink) - (Context, Devices.size(), Devices.data(), - Opts, LinkPrograms.size(), &LinkPrograms[0], - nullptr, nullptr, &LinkedProg); + pi_result Error = PI_CALL_NOCHECK(piProgramLink)( + Context, Devices.size(), Devices.data(), Opts, LinkPrograms.size(), + &LinkPrograms[0], nullptr, nullptr, &LinkedProg); if (Error != PI_SUCCESS) { throw compile_program_error(getProgramBuildLog(Program.get())); } diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp index 2c780fee8b768..d8539db531ec8 100644 --- a/sycl/test/devicelib/assert-windows.cpp +++ b/sycl/test/devicelib/assert-windows.cpp @@ -16,9 +16,9 @@ // RUN: FileCheck %s --input-file %t.stdout.fallback --check-prefix=CHECK-FALLBACK // CHECK-FALLBACK: ---> piProgramLink -#include -#include #include +#include +#include using namespace cl::sycl; @@ -26,35 +26,34 @@ constexpr auto sycl_read = cl::sycl::access::mode::read; constexpr auto sycl_write = cl::sycl::access::mode::write; template -void simple_vadd(const std::array& VA, const std::array& VB, - std::array& VC) { +void simple_vadd(const std::array &VA, const std::array &VB, + std::array &VC) { queue deviceQueue([](cl::sycl::exception_list ExceptionList) { - for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { - try { - std::rethrow_exception(ExceptionPtr); - } catch (cl::sycl::exception &E) { - std::cerr << E.what() << std::endl; - } catch (...) { - std::cerr << "Unknown async exception was caught." << std::endl; - } - } - }); + for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { + try { + std::rethrow_exception(ExceptionPtr); + } catch (cl::sycl::exception &E) { + std::cerr << E.what() << std::endl; + } catch (...) { + std::cerr << "Unknown async exception was caught." << std::endl; + } + } + }); cl::sycl::range<1> numOfItems{N}; cl::sycl::buffer bufferA(VA.data(), numOfItems); cl::sycl::buffer bufferB(VB.data(), numOfItems); cl::sycl::buffer bufferC(VC.data(), numOfItems); - deviceQueue.submit([&](cl::sycl::handler& cgh) { + deviceQueue.submit([&](cl::sycl::handler &cgh) { auto accessorA = bufferA.template get_access(cgh); auto accessorB = bufferB.template get_access(cgh); auto accessorC = bufferC.template get_access(cgh); - cgh.parallel_for(numOfItems, - [=](cl::sycl::id<1> wiID) { - accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; - assert(accessorC[wiID] == 0 && "Invalid value"); - }); + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + assert(accessorC[wiID] == 0 && "Invalid value"); + }); }); deviceQueue.wait_and_throw(); } diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp index 2b4c69d9d252f..97abee4b63315 100644 --- a/sycl/test/devicelib/assert.cpp +++ b/sycl/test/devicelib/assert.cpp @@ -92,14 +92,13 @@ // Note that the work-item that hits the assert first may vary, since the order // of execution is undefined. We catch only the first one (whatever id it is). -#include -#include #include - +#include +#include +#include #include #include #include -#include using namespace cl::sycl; @@ -109,19 +108,19 @@ constexpr auto sycl_write = cl::sycl::access::mode::write; const int EXIT_SKIP_TEST = 42; template -void simple_vadd(const std::array& VA, const std::array& VB, - std::array& VC) { +void simple_vadd(const std::array &VA, const std::array &VB, + std::array &VC) { queue deviceQueue([](cl::sycl::exception_list ExceptionList) { - for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { - try { - std::rethrow_exception(ExceptionPtr); - } catch (cl::sycl::exception &E) { - std::cerr << E.what() << std::endl; - } catch (...) { - std::cerr << "Unknown async exception was caught." << std::endl; - } - } - }); + for (cl::sycl::exception_ptr_class ExceptionPtr : ExceptionList) { + try { + std::rethrow_exception(ExceptionPtr); + } catch (cl::sycl::exception &E) { + std::cerr << E.what() << std::endl; + } catch (...) { + std::cerr << "Unknown async exception was caught." << std::endl; + } + } + }); device dev = deviceQueue.get_device(); bool unsupported = true; for (auto &ext : dev.get_info()) { @@ -131,7 +130,7 @@ void simple_vadd(const std::array& VA, const std::array& VB, } if (unsupported && getenv("SKIP_IF_NO_EXT")) { fprintf(stderr, "Device has no support for cl_intel_devicelib_assert, " - "skipping the test\n"); + "skipping the test\n"); exit(EXIT_SKIP_TEST); } @@ -140,16 +139,15 @@ void simple_vadd(const std::array& VA, const std::array& VB, cl::sycl::buffer bufferB(VB.data(), numOfItems); cl::sycl::buffer bufferC(VC.data(), numOfItems); - deviceQueue.submit([&](cl::sycl::handler& cgh) { + deviceQueue.submit([&](cl::sycl::handler &cgh) { auto accessorA = bufferA.template get_access(cgh); auto accessorB = bufferB.template get_access(cgh); auto accessorC = bufferC.template get_access(cgh); - cgh.parallel_for(numOfItems, - [=](cl::sycl::id<1> wiID) { - accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; - assert(accessorC[wiID] == 0 && "Invalid value"); - }); + cgh.parallel_for(numOfItems, [=](cl::sycl::id<1> wiID) { + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + assert(accessorC[wiID] == 0 && "Invalid value"); + }); }); deviceQueue.wait_and_throw(); } @@ -168,16 +166,15 @@ int main() { } int sig = WTERMSIG(status); int expected = 0; - if (const char* env = getenv("EXPECTED_SIGNAL")) { + if (const char *env = getenv("EXPECTED_SIGNAL")) { if (0 == strcmp(env, "SIGABRT")) { expected = SIGABRT; } else if (0 == strcmp(env, "SIGSEGV")) { expected = SIGSEGV; } if (!expected) { - fprintf(stderr, - "EXPECTED_SIGNAL should be set to either \"SIGABRT\", " - "or \"SIGSEGV\"!\n"); + fprintf(stderr, "EXPECTED_SIGNAL should be set to either \"SIGABRT\", " + "or \"SIGSEGV\"!\n"); return 1; } } From 7b57b6ba310b84e1bea9227114656dff1d6c9d8e Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 19:58:33 +0300 Subject: [PATCH 06/32] Code style change Signed-off-by: Andrew Savonichev --- .../detail/program_manager/program_manager.cpp | 15 ++++++++------- 1 file changed, 8 insertions(+), 7 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index eade277cf5f57..47eca055322fc 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -310,12 +310,13 @@ loadDeviceLibFallback(const RT::PiContext &Context, Extension); } RT::PiProgram &LibProg = CachedLibPrograms[LibFileName]; - bool ShouldCompile = !LibProg; - if (!LibProg && !loadDeviceLib(Context, LibFileName, LibProg)) + if (LibProg) + return LibProg; + + if (!loadDeviceLib(Context, LibFileName, LibProg)) throw compile_program_error(std::string("Failed to load ") + LibFileName); - if (ShouldCompile) - PI_CALL(piProgramCompile) + PI_CALL(piProgramCompile) (LibProg, // Assume that Devices contains all devices from Context. Devices.size(), Devices.data(), @@ -485,9 +486,9 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, if (LinkPrograms.empty()) { pi_result Error = PI_CALL_NOCHECK(piProgramBuild)( Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr); - if (Error == PI_SUCCESS) - return Program; - throw compile_program_error(getProgramBuildLog(Program.get())); + if (Error != PI_SUCCESS) + compile_program_error(getProgramBuildLog(Program.get())); + return Program; } // Include the main program and compile/link everything together From 462af28b626a392e860066fdf4602772c14b2c2c Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 21:26:27 +0300 Subject: [PATCH 07/32] Release library programs Signed-off-by: Andrew Savonichev --- sycl/source/detail/context_impl.cpp | 3 +++ 1 file changed, 3 insertions(+) diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index 7a2f8218ebdc9..dcafd3ea4ea0f 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -90,6 +90,9 @@ context_impl::~context_impl() { PI_CALL(piKernelRelease)(KernIt.second); PI_CALL(piProgramRelease)(ToBeDeleted); } + for (auto LibProg : MCachedLibPrograms) { + PI_CALL(piProgramRelease)(LibProg.second); + } } const async_handler &context_impl::get_async_handler() const { From 52e077c8e5e0ebe5905695aeaba869a1ef6c1e2a Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 21:27:13 +0300 Subject: [PATCH 08/32] Cleanup invalid library programs if an exception is thrown Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.cpp | 20 +++++++++++++++---- 1 file changed, 16 insertions(+), 4 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 47eca055322fc..1630ddc0b3bb0 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -309,14 +309,22 @@ loadDeviceLibFallback(const RT::PiContext &Context, throw compile_program_error(std::string("Unknown device library: ") + Extension); } - RT::PiProgram &LibProg = CachedLibPrograms[LibFileName]; - if (LibProg) + std::map::iterator LibProgIt; + bool NotExists = false; + std::tie(LibProgIt, NotExists) = + CachedLibPrograms.insert({Extension, nullptr}); + RT::PiProgram &LibProg = LibProgIt->second; + + if (!NotExists) { return LibProg; + } - if (!loadDeviceLib(Context, LibFileName, LibProg)) + if (!loadDeviceLib(Context, LibFileName, LibProg)) { + CachedLibPrograms.erase(LibProgIt); throw compile_program_error(std::string("Failed to load ") + LibFileName); + } - PI_CALL(piProgramCompile) + pi_result Error = PI_CALL_NOCHECK(piProgramCompile) (LibProg, // Assume that Devices contains all devices from Context. Devices.size(), Devices.data(), @@ -325,6 +333,10 @@ loadDeviceLibFallback(const RT::PiContext &Context, // library program as well, and what actually happens to a SPIR-V // program if we apply them. "", 0, nullptr, nullptr, nullptr, nullptr); + if (Error != PI_SUCCESS) { + CachedLibPrograms.erase(LibProgIt); + throw compile_program_error(ProgramManager::getProgramBuildLog(LibProg)); + } return LibProg; } From 75c05b589ff276f8c25565f6d754d741ff8605a1 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 21:29:51 +0300 Subject: [PATCH 09/32] Tabs vs spaces Signed-off-by: Andrew Savonichev --- .../C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst | 6 +++--- .../C-CXX-StandardLibrary/DeviceLibExtensions.rst | 2 +- 2 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst index 30a2d156590e8..e3d26f808a53f 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst @@ -35,7 +35,7 @@ Example of usage template void simple_vadd(const std::array& VA, const std::array& VB, - std::array& VC) { + std::array& VC) { // ... cl::sycl::range<1> numOfItems{N}; cl::sycl::buffer bufferA(VA.data(), numOfItems); @@ -49,8 +49,8 @@ Example of usage cgh.parallel_for>(numOfItems, [=](cl::sycl::id<1> wiID) { - accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; - assert(accessorC[wiID] > 0 && "Invalid value"); + accessorC[wiID] = accessorA[wiID] + accessorB[wiID]; + assert(accessorC[wiID] > 0 && "Invalid value"); }); }); deviceQueue.wait_and_throw(); diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index 6596dd5b16ac9..c4295188dc01c 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -11,7 +11,7 @@ cl_intel_devicelib_cassert void __devicelib_assert_fail(__generic const char *expr, __generic const char *file, int32_t line, - __generic const char *func, + __generic const char *func, size_t gid0, size_t gid1, size_t gid2, size_t lid0, size_t lid1, size_t lid2); Semantic: From 269d6dc4a043a3ff7ec9b46c43988993f8fb384c Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 11 Dec 2019 21:32:50 +0300 Subject: [PATCH 10/32] Reword documentation in context_impl.hpp Signed-off-by: Andrew Savonichev --- sycl/include/CL/sycl/detail/context_impl.hpp | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl/include/CL/sycl/detail/context_impl.hpp b/sycl/include/CL/sycl/detail/context_impl.hpp index be4803a30388f..2b99bb8aacc19 100644 --- a/sycl/include/CL/sycl/detail/context_impl.hpp +++ b/sycl/include/CL/sycl/detail/context_impl.hpp @@ -122,16 +122,18 @@ class context_impl { /// @return a pointer to USM dispatcher. std::shared_ptr getUSMDispatch() const; - /// Returns device library programs: in contrast to user programs, which are a - /// part of user code, these programs come from the SYCL runtime. Device - /// libraries are identified by the corresponding extension name: + /// 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 name: /// /// "cl_intel_devicelib_assert" -> # /// "cl_intel_devicelib_complex" -> # - //// etc. + /// etc. /// /// See `doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst' for /// more details. + /// + /// @returns a map with device library programs. std::map &getCachedLibPrograms() { return MCachedLibPrograms; } From 3d80fe0541a2459c43607c4b45383f7b75bbab6d Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 12 Dec 2019 12:38:32 +0300 Subject: [PATCH 11/32] Fix error handling of clLinkProgram Signed-off-by: Andrew Savonichev --- .../detail/program_manager/program_manager.cpp | 14 +++++++++++--- 1 file changed, 11 insertions(+), 3 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1630ddc0b3bb0..1f41ed302f370 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -513,11 +513,19 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, pi_result Error = PI_CALL_NOCHECK(piProgramLink)( Context, Devices.size(), Devices.data(), Opts, LinkPrograms.size(), &LinkPrograms[0], nullptr, nullptr, &LinkedProg); + + // Link program call returns a new program object if all parameters are valid, + // or NULL otherwise. Release the original (user) program. + Program.reset(LinkedProg); if (Error != PI_SUCCESS) { - throw compile_program_error(getProgramBuildLog(Program.get())); + if (LinkedProg) { + // A non-trivial error occurred during linkage: get a build log, release + // an incomplete (but valid) LinkedProg, and throw. + throw compile_program_error(getProgramBuildLog(LinkedProg)); + } + pi::checkPiResult(Error); } - return ProgramPtr(LinkedProg, - RT::PluginInformation.PiFunctionTable.piProgramRelease); + return Program; } void ProgramManager::addImages(pi_device_binaries DeviceBinary) { From 7497a23af61ac61a0c86b2e5190ac00e302b4f92 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 12 Dec 2019 13:59:20 +0300 Subject: [PATCH 12/32] Code style changes Signed-off-by: Andrew Savonichev --- sycl/source/detail/program_manager/program_manager.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 1f41ed302f370..eea6ad77d38a1 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -324,7 +324,7 @@ loadDeviceLibFallback(const RT::PiContext &Context, throw compile_program_error(std::string("Failed to load ") + LibFileName); } - pi_result Error = PI_CALL_NOCHECK(piProgramCompile) + RT::PiResult Error = PI_CALL_NOCHECK(piProgramCompile) (LibProg, // Assume that Devices contains all devices from Context. Devices.size(), Devices.data(), @@ -438,6 +438,7 @@ void getDeviceLibPrograms( const std::vector Devices, std::map &CachedLibPrograms, std::vector &Programs) { + // TODO: SYCL compiler should generate a list of required extensions for a // particular program in order to allow us do a more fine-grained check here. // Require *all* possible devicelib extensions for now. @@ -455,6 +456,8 @@ void getDeviceLibPrograms( InhibitNativeImpl = strstr(Env, Ext) != nullptr; } + // Load a fallback library for an extension if at least one device does not + // support it. for (const std::string &DevExtList : DevExtensions) { bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); if (!DeviceSupports || InhibitNativeImpl) { @@ -496,7 +499,7 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, } if (LinkPrograms.empty()) { - pi_result Error = PI_CALL_NOCHECK(piProgramBuild)( + RT::PiResult Error = PI_CALL_NOCHECK(piProgramBuild)( Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr); if (Error != PI_SUCCESS) compile_program_error(getProgramBuildLog(Program.get())); @@ -510,7 +513,7 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, LinkPrograms.push_back(Program.get()); RT::PiProgram LinkedProg = nullptr; - pi_result Error = PI_CALL_NOCHECK(piProgramLink)( + RT::PiResult Error = PI_CALL_NOCHECK(piProgramLink)( Context, Devices.size(), Devices.data(), Opts, LinkPrograms.size(), &LinkPrograms[0], nullptr, nullptr, &LinkedProg); From 82e0e8485d8318e1c42e92b823f6c28b7c9742af Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 12 Dec 2019 14:18:32 +0300 Subject: [PATCH 13/32] Include path to SYCL headers This should fix: fatal error: 'CL/__spirv/spirv_vars.hpp' file not found Signed-off-by: Andrew Savonichev --- sycl/source/detail/devicelib/CMakeLists.txt | 15 +++++++++------ 1 file changed, 9 insertions(+), 6 deletions(-) diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt index 92966b2566698..28f9b25311025 100644 --- a/sycl/source/detail/devicelib/CMakeLists.txt +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -8,12 +8,17 @@ endif() set(clang $) +set(compile_opts + # suppress an error about SYCL_EXTERNAL + -Wno-error=sycl-strict -Wno-sycl-strict + # for CL/__spirv/spirv_vars.hpp + -I${sycl_inc_dir}) + if (MSVC) set(devicelib-obj-file ${binary_dir}/libsycl-msvc.o) add_custom_command(OUTPUT ${devicelib-obj-file} COMMAND ${clang} -fsycl -c - # suppress an error about SYCL_EXTERNAL - -Wno-error=sycl-strict -Wno-sycl-strict + ${compile_opts} ${CMAKE_CURRENT_SOURCE_DIR}/msvc_wrapper.cpp -o ${devicelib-obj-file} MAIN_DEPENDENCY msvc_wrapper.cpp @@ -23,8 +28,7 @@ else() set(devicelib-obj-file ${binary_dir}/libsycl-glibc.o) add_custom_command(OUTPUT ${devicelib-obj-file} COMMAND ${clang} -fsycl -c - # suppress an error about SYCL_EXTERNAL - -Wno-error=sycl-strict -Wno-sycl-strict + ${compile_opts} ${CMAKE_CURRENT_SOURCE_DIR}/glibc_wrapper.cpp -o ${devicelib-obj-file} MAIN_DEPENDENCY glibc_wrapper.cpp @@ -34,8 +38,7 @@ endif() add_custom_command(OUTPUT ${binary_dir}/libsycl-fallback-cassert.spv COMMAND ${clang} -fsycl-device-only -S -Xclang -emit-llvm-bc - # suppress an error about SYCL_EXTERNAL - -Wno-error=sycl-strict -Wno-sycl-strict + ${compile_opts} ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp -o libsycl-fallback-cassert.bc COMMAND llvm-spirv libsycl-fallback-cassert.bc From 265aaf6e5f5bc2bad8d911d7c03e0a7730b472d7 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 13 Dec 2019 17:07:48 +0300 Subject: [PATCH 14/32] Re-format PI_CALL Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.cpp | 47 +++++++++---------- 1 file changed, 22 insertions(+), 25 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index eea6ad77d38a1..a179e7977d4a0 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -91,13 +91,12 @@ static RT::PiProgram createSpirvProgram(const RT::PiContext Context, static void getContextDevices(const RT::PiContext &Context, std::vector &Devices) { size_t NumDevices = 0; - PI_CALL(piContextGetInfo) - (Context, PI_CONTEXT_INFO_NUM_DEVICES, sizeof(NumDevices), &NumDevices, - nullptr); + PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_NUM_DEVICES, + sizeof(NumDevices), &NumDevices, nullptr); Devices.resize(NumDevices); - PI_CALL(piContextGetInfo) - (Context, PI_CONTEXT_INFO_DEVICES, sizeof(RT::PiDevice) * Devices.size(), - &Devices[0], nullptr); + PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_DEVICES, + sizeof(RT::PiDevice) * Devices.size(), &Devices[0], + nullptr); } DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, @@ -285,14 +284,13 @@ static bool loadDeviceLib(const RT::PiContext &Context, const char *Name, static std::string getDeviceExtensions(const RT::PiDevice &Dev) { std::string DevExt; size_t DevExtSize = 0; - PI_CALL(piDeviceGetInfo) - (Dev, PI_DEVICE_INFO_EXTENSIONS, - /*param_value_size=*/0, - /*param_value=*/nullptr, &DevExtSize); + PI_CALL(piDeviceGetInfo)(Dev, PI_DEVICE_INFO_EXTENSIONS, + /*param_value_size=*/0, + /*param_value=*/nullptr, &DevExtSize); DevExt.resize(DevExtSize); - PI_CALL(piDeviceGetInfo) - (Dev, PI_DEVICE_INFO_EXTENSIONS, DevExt.size(), &DevExt[0], - /*param_value_size_ret=*/nullptr); + PI_CALL(piDeviceGetInfo)(Dev, PI_DEVICE_INFO_EXTENSIONS, DevExt.size(), + &DevExt[0], + /*param_value_size_ret=*/nullptr); return DevExt; } @@ -324,15 +322,15 @@ loadDeviceLibFallback(const RT::PiContext &Context, throw compile_program_error(std::string("Failed to load ") + LibFileName); } - RT::PiResult Error = PI_CALL_NOCHECK(piProgramCompile) - (LibProg, - // Assume that Devices contains all devices from Context. - Devices.size(), Devices.data(), - // Do not use compile options for library programs: it is not clear - // if user options (image options) are supposed to be applied to - // library program as well, and what actually happens to a SPIR-V - // program if we apply them. - "", 0, nullptr, nullptr, nullptr, nullptr); + RT::PiResult Error = PI_CALL_NOCHECK(piProgramCompile)( + LibProg, + // Assume that Devices contains all devices from Context. + Devices.size(), Devices.data(), + // Do not use compile options for library programs: it is not clear + // if user options (image options) are supposed to be applied to + // library program as well, and what actually happens to a SPIR-V + // program if we apply them. + "", 0, nullptr, nullptr, nullptr, nullptr); if (Error != PI_SUCCESS) { CachedLibPrograms.erase(LibProgIt); throw compile_program_error(ProgramManager::getProgramBuildLog(LibProg)); @@ -507,9 +505,8 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, } // Include the main program and compile/link everything together - PI_CALL(piProgramCompile) - (Program.get(), Devices.size(), Devices.data(), Opts, 0, nullptr, nullptr, - nullptr, nullptr); + PI_CALL(piProgramCompile)(Program.get(), Devices.size(), Devices.data(), Opts, + 0, nullptr, nullptr, nullptr, nullptr); LinkPrograms.push_back(Program.get()); RT::PiProgram LinkedProg = nullptr; From 4ce3a5760a17dc838f638a44a223a662a25e0364 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 13 Dec 2019 17:14:05 +0300 Subject: [PATCH 15/32] Misc code review comments Signed-off-by: Andrew Savonichev --- .../C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst | 2 +- .../CL/sycl/detail/program_manager/program_manager.hpp | 2 +- sycl/source/detail/context_impl.cpp | 1 + sycl/source/detail/devicelib/CMakeLists.txt | 6 ++---- sycl/source/detail/program_manager/program_manager.cpp | 2 +- 5 files changed, 6 insertions(+), 7 deletions(-) diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst index e3d26f808a53f..69f14ff73288e 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst @@ -113,7 +113,7 @@ For example, `__assert_fail` from IR above gets transformed into: A single wrapper object provides function wrappers for *all* supported library functions. Every supported C library implementation (MSVC or -glibc) have its own wrapper library object: +glibc) has its own wrapper library object: - libsycl-glibc.o - libsycl-msvc.o diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index 802562ebf81d8..ef6b0a25aa2e3 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp @@ -74,7 +74,7 @@ class ProgramManager { decltype(&::piProgramRelease)>; ProgramPtr build(ProgramPtr Program, RT::PiContext Context, const string_class &Options, - std::vector Devices, + const std::vector &Devices, std::map &CachedLibPrograms, bool LinkDeviceLibs = false); /// Provides a new kernel set id for grouping kernel names together diff --git a/sycl/source/detail/context_impl.cpp b/sycl/source/detail/context_impl.cpp index dcafd3ea4ea0f..d0520b9d56b01 100644 --- a/sycl/source/detail/context_impl.cpp +++ b/sycl/source/detail/context_impl.cpp @@ -91,6 +91,7 @@ context_impl::~context_impl() { 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); } } diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt index 28f9b25311025..1db54f605263d 100644 --- a/sycl/source/detail/devicelib/CMakeLists.txt +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -10,7 +10,7 @@ set(clang $) set(compile_opts # suppress an error about SYCL_EXTERNAL - -Wno-error=sycl-strict -Wno-sycl-strict + -Wno-sycl-strict # for CL/__spirv/spirv_vars.hpp -I${sycl_inc_dir}) @@ -37,11 +37,9 @@ else() endif() add_custom_command(OUTPUT ${binary_dir}/libsycl-fallback-cassert.spv - COMMAND ${clang} -fsycl-device-only -S -Xclang -emit-llvm-bc + COMMAND ${clang} -S -fsycl-device-only -fno-sycl-use-bitcode ${compile_opts} ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp - -o libsycl-fallback-cassert.bc - COMMAND llvm-spirv libsycl-fallback-cassert.bc -o ${binary_dir}/libsycl-fallback-cassert.spv MAIN_DEPENDENCY fallback-cassert.cpp DEPENDS wrapper.h clang diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index a179e7977d4a0..9be9e08b9529a 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -470,7 +470,7 @@ void getDeviceLibPrograms( ProgramManager::ProgramPtr ProgramManager::build(ProgramPtr Program, RT::PiContext Context, const string_class &Options, - std::vector Devices, + const std::vector &Devices, std::map &CachedLibPrograms, bool LinkDeviceLibs) { From 278206f0caf3f7ee9b9249f381bb7390b664a9e9 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 13 Dec 2019 17:18:52 +0300 Subject: [PATCH 16/32] Misc code review comments Signed-off-by: Andrew Savonichev --- sycl/source/detail/program_manager/program_manager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 9be9e08b9529a..aeacba3e052b4 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -433,7 +433,7 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, void getDeviceLibPrograms( const RT::PiContext Context, - const std::vector Devices, + const std::vector &Devices, std::map &CachedLibPrograms, std::vector &Programs) { From a5d65dbb7bb069f7059e7a34678d00df3de210d4 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Fri, 13 Dec 2019 20:34:23 +0300 Subject: [PATCH 17/32] Add a dependency for llvm-spirv Clang calls it under the hood, but it has no explicit dependency for it in cmake. Signed-off-by: Andrew Savonichev --- sycl/source/detail/devicelib/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt index 1db54f605263d..c6a7e28efbc99 100644 --- a/sycl/source/detail/devicelib/CMakeLists.txt +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -42,7 +42,7 @@ add_custom_command(OUTPUT ${binary_dir}/libsycl-fallback-cassert.spv ${CMAKE_CURRENT_SOURCE_DIR}/fallback-cassert.cpp -o ${binary_dir}/libsycl-fallback-cassert.spv MAIN_DEPENDENCY fallback-cassert.cpp - DEPENDS wrapper.h clang + DEPENDS wrapper.h clang llvm-spirv VERBATIM) add_custom_target(devicelib-obj DEPENDS ${devicelib-obj-file}) From 6b676010fee2062210534aa2c292088c8ede532f Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Mon, 16 Dec 2019 17:11:41 +0300 Subject: [PATCH 18/32] Disable the test for assert on Windows Signed-off-by: Andrew Savonichev --- sycl/test/devicelib/assert-windows.cpp | 5 +++++ 1 file changed, 5 insertions(+) diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp index d8539db531ec8..ed0970038d3f2 100644 --- a/sycl/test/devicelib/assert-windows.cpp +++ b/sycl/test/devicelib/assert-windows.cpp @@ -1,4 +1,9 @@ // REQUIRES: cpu,windows +// +// FIXME: OpenCL CPU backend compiler crashes on a call to _wassert. +// Disable the test until the fix reaches SYCL test infrastructure. +// XFAIL: * +// // RUN: %clangxx -fsycl -c %s -o %t.o // RUN: %clangxx -fsycl %t.o %llvm_build_libs_dir/../bin/libsycl-msvc.o -o %t.out // From d5859219ad20aaf7b5d63458fd30f0a614428d48 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 17 Dec 2019 18:01:01 +0300 Subject: [PATCH 19/32] Reword documentation Signed-off-by: Andrew Savonichev --- .../C-CXX-StandardLibrary.rst | 32 ++++++++++--------- .../DeviceLibExtensions.rst | 8 ++--- 2 files changed, 21 insertions(+), 19 deletions(-) diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst index 69f14ff73288e..42e3c22f52664 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst @@ -1,22 +1,18 @@ C and C++ Standard libraries support =================================== -This extension enables a set of functions from 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 -or ), and the corresponding header has to be explicitly -included in user code. +This extension enables a set of functions from the C and C++ standard +libraries, and allows to use them in SYCL device code. -List of supported functions from C standard library: - - assert macro (from assert.h) - -NOTE: only the GNU glibc and Microsoft C libraries are currently -supported. +Function declarations are taken from the standard headers (e.g. from + or ), and the corresponding header has to be +explicitly included in user code. -Device library is distributed with the compiler, and it has to be -explicitly linked by a user. +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: -On Linux with GNU glibc: +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 @@ -26,6 +22,12 @@ or, in case of Windows: 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) + +NOTE: only the GNU glibc and Microsoft C libraries are currently +supported. + Example of usage ================ @@ -72,7 +74,7 @@ functions in LLVM IR after clang: call spir_func void @__assert_fail([...]) unreachable -C and C++ specifications do not define names and signatures of the +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: @@ -82,7 +84,7 @@ function. For example, the `assert` macro: 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, a +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. diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index c4295188dc01c..c368ba5d60bd0 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -17,15 +17,15 @@ cl_intel_devicelib_cassert 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 -the arguments. In addition to that, the function is free to terminate -the current kernel invocation. +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 a function containing 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 From e8824600b4869a6f7ad88446018b1c71d0f9412c Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 17 Dec 2019 18:01:17 +0300 Subject: [PATCH 20/32] Use 'if WIN32' instead of 'if MSVC' in CMake Signed-off-by: Andrew Savonichev --- sycl/source/detail/devicelib/CMakeLists.txt | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/source/detail/devicelib/CMakeLists.txt b/sycl/source/detail/devicelib/CMakeLists.txt index c6a7e28efbc99..273110a783496 100644 --- a/sycl/source/detail/devicelib/CMakeLists.txt +++ b/sycl/source/detail/devicelib/CMakeLists.txt @@ -1,6 +1,6 @@ # Place device libraries near the libsycl.so library in a build # directory -if (MSVC) +if (WIN32) set(binary_dir "${CMAKE_RUNTIME_OUTPUT_DIRECTORY}") else() set(binary_dir "${CMAKE_LIBRARY_OUTPUT_DIRECTORY}") @@ -14,7 +14,7 @@ set(compile_opts # for CL/__spirv/spirv_vars.hpp -I${sycl_inc_dir}) -if (MSVC) +if (WIN32) set(devicelib-obj-file ${binary_dir}/libsycl-msvc.o) add_custom_command(OUTPUT ${devicelib-obj-file} COMMAND ${clang} -fsycl -c @@ -54,7 +54,7 @@ endif() # Place device libraries near the libsycl.so library in an install # directory as well -if (MSVC) +if (WIN32) set(install_dest bin) else() set(install_dest lib) From f1ae75f4623a353204c283ff04b2e58a5382d3b4 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Tue, 17 Dec 2019 21:18:51 +0300 Subject: [PATCH 21/32] Code review comments Signed-off-by: Andrew Savonichev --- sycl/source/detail/program_manager/program_manager.cpp | 8 ++++---- sycl/test/devicelib/assert-windows.cpp | 2 +- sycl/test/devicelib/assert.cpp | 2 +- 3 files changed, 6 insertions(+), 6 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index aeacba3e052b4..aa605f4524170 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -194,7 +194,7 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, // Link a fallback implementation of device libraries if they are not // supported by a device compiler. // Pre-compiled programs are supposed to be already linked. - bool LinkDeviceLibs = getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; + const bool LinkDeviceLibs = getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; std::vector Devices; getContextDevices(getRawSyclObjImpl(Context)->getHandleRef(), Devices); @@ -431,7 +431,7 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, return *Img; } -void getDeviceLibPrograms( +static void getDeviceLibPrograms( const RT::PiContext Context, const std::vector &Devices, std::map &CachedLibPrograms, @@ -500,7 +500,7 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, RT::PiResult Error = PI_CALL_NOCHECK(piProgramBuild)( Program.get(), Devices.size(), Devices.data(), Opts, nullptr, nullptr); if (Error != PI_SUCCESS) - compile_program_error(getProgramBuildLog(Program.get())); + throw compile_program_error(getProgramBuildLog(Program.get())); return Program; } @@ -512,7 +512,7 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, RT::PiProgram LinkedProg = nullptr; RT::PiResult Error = PI_CALL_NOCHECK(piProgramLink)( Context, Devices.size(), Devices.data(), Opts, LinkPrograms.size(), - &LinkPrograms[0], nullptr, nullptr, &LinkedProg); + LinkPrograms.data(), nullptr, nullptr, &LinkedProg); // Link program call returns a new program object if all parameters are valid, // or NULL otherwise. Release the original (user) program. diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp index ed0970038d3f2..fe2700a1b5382 100644 --- a/sycl/test/devicelib/assert-windows.cpp +++ b/sycl/test/devicelib/assert-windows.cpp @@ -14,7 +14,7 @@ // explicitly. Since the test is going to crash, we'll have to follow a similar // approach as on Linux - call the test in a subprocess. // -// RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_LINK_FALLBACK=1 CL_CONFIG_USE_VECTORIZER=False SYCL_DEVICE_TYPE=CPU %t.out >%t.stdout.fallback 2>%t.stderr.fallback +// RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_INHIBIT_NATIVE=1 SYCL_DEVICE_TYPE=CPU %t.out >%t.stdout.fallback 2>%t.stderr.fallback // RUN: FileCheck %s --check-prefix=CHECK-MESSAGE --input-file %t.stdout.fallback // CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: : local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. // diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp index 97abee4b63315..aa6b9702b1f1e 100644 --- a/sycl/test/devicelib/assert.cpp +++ b/sycl/test/devicelib/assert.cpp @@ -4,7 +4,7 @@ // (see the other RUN lines below; it is a bit complicated) // // assert() call in device code guarantees nothing: on some devices it behaves -// in a usual way and terminate a program. On other devices it can print an +// in the usual way and terminate the program. On other devices it can print an // error message and *continue* execution. Less capable devices can even ignore // an assert! // From a1aeda7985c8f41b47fb7dde8c10b7345b777a47 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 13:13:18 +0300 Subject: [PATCH 22/32] Get devices from context instead of going through OpenCL API Signed-off-by: Andrew Savonichev --- sycl/include/CL/sycl/detail/context_impl.hpp | 6 ++++ sycl/include/CL/sycl/device.hpp | 6 ++++ .../program_manager/program_manager.cpp | 29 +++++++------------ 3 files changed, 23 insertions(+), 18 deletions(-) diff --git a/sycl/include/CL/sycl/detail/context_impl.hpp b/sycl/include/CL/sycl/detail/context_impl.hpp index 2b99bb8aacc19..983b3abffe681 100644 --- a/sycl/include/CL/sycl/detail/context_impl.hpp +++ b/sycl/include/CL/sycl/detail/context_impl.hpp @@ -102,6 +102,12 @@ class context_impl { /// @return an instance of raw plug-in context handle. const RT::PiContext &getHandleRef() const; + /// Unlike `get_info', this function returns a + /// reference. + const vector_class &getDevices() const { + return MDevices; + } + /// Gets cached programs. /// /// @return a map of cached programs. diff --git a/sycl/include/CL/sycl/device.hpp b/sycl/include/CL/sycl/device.hpp index fb14f19ecc772..eac87b89880fa 100644 --- a/sycl/include/CL/sycl/device.hpp +++ b/sycl/include/CL/sycl/device.hpp @@ -157,6 +157,12 @@ class device { template friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); + + template + friend + typename std::add_pointer::type + detail::getRawSyclObjImpl(const T &SyclObject); + template friend T detail::createSyclObjFromImpl(decltype(T::impl) ImplObj); }; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index aa605f4524170..626c68caac3c6 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -88,17 +88,6 @@ static RT::PiProgram createSpirvProgram(const RT::PiContext Context, return Program; } -static void getContextDevices(const RT::PiContext &Context, - std::vector &Devices) { - size_t NumDevices = 0; - PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_NUM_DEVICES, - sizeof(NumDevices), &NumDevices, nullptr); - Devices.resize(NumDevices); - PI_CALL(piContextGetInfo)(Context, PI_CONTEXT_INFO_DEVICES, - sizeof(RT::PiDevice) * Devices.size(), &Devices[0], - nullptr); -} - DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, const string_class &KernelName, const context &Context) { @@ -196,13 +185,17 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, // Pre-compiled programs are supposed to be already linked. const bool LinkDeviceLibs = getFormat(Img) == PI_DEVICE_BINARY_TYPE_SPIRV; - std::vector Devices; - getContextDevices(getRawSyclObjImpl(Context)->getHandleRef(), Devices); - - ProgramPtr BuiltProgram = build( - std::move(ProgramManaged), getRawSyclObjImpl(Context)->getHandleRef(), - Img.BuildOptions, Devices, - getRawSyclObjImpl(Context)->getCachedLibPrograms(), LinkDeviceLibs); + context_impl *ContextImpl = getRawSyclObjImpl(Context); + RT::PiContext PiContext = ContextImpl->getHandleRef(); + const std::vector Devices = ContextImpl->getDevices(); + std::vector PiDevices(Devices.size()); + std::transform( + Devices.begin(), Devices.end(), PiDevices.begin(), + [](const device Dev) { return getRawSyclObjImpl(Dev)->getHandleRef(); }); + + ProgramPtr BuiltProgram = + build(std::move(ProgramManaged), PiContext, Img.BuildOptions, PiDevices, + ContextImpl->getCachedLibPrograms(), LinkDeviceLibs); CachedPrograms[KSId] = BuiltProgram.get(); return BuiltProgram.release(); } From f3a4e0747b2ccd451309e35d318b6cb9b66427c0 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 14:11:45 +0300 Subject: [PATCH 23/32] Refactor getDeviceLibPrograms Signed-off-by: Andrew Savonichev --- .../program_manager/program_manager.cpp | 45 +++++++++++-------- 1 file changed, 26 insertions(+), 19 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 626c68caac3c6..d9ec59df49201 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -424,40 +424,47 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, return *Img; } -static void getDeviceLibPrograms( +static std::vector getDeviceLibPrograms( const RT::PiContext Context, const std::vector &Devices, - std::map &CachedLibPrograms, - std::vector &Programs) { + std::map &CachedLibPrograms) { + + std::vector Programs; // TODO: SYCL compiler should generate a list of required extensions for a // particular program in order to allow us do a more fine-grained check here. // Require *all* possible devicelib extensions for now. - const char* RequiredDeviceLibExt[] = { - "cl_intel_devicelib_assert" + std::pair RequiredDeviceLibExt[] = { + {"cl_intel_devicelib_assert", false} }; - std::vector DevExtensions(Devices.size()); - for (size_t i = 0; i < Devices.size(); ++i) { - DevExtensions[i] = getDeviceExtensions(Devices[i]); - } - for (const char *Ext : RequiredDeviceLibExt) { - bool InhibitNativeImpl = false; - if (const char *Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { - InhibitNativeImpl = strstr(Env, Ext) != nullptr; - } + // Load a fallback library for an extension if at least one device does not + // support it. + for (RT::PiDevice Dev : Devices) { + std::string DevExtList = getDeviceExtensions(Dev); + for (auto &Pair : RequiredDeviceLibExt) { + const char *Ext = Pair.first; + bool &FallbackIsLoaded = Pair.second; + + if (FallbackIsLoaded) { + continue; + } + + bool InhibitNativeImpl = false; + if (const char *Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { + InhibitNativeImpl = strstr(Env, Ext) != nullptr; + } - // Load a fallback library for an extension if at least one device does not - // support it. - for (const std::string &DevExtList : DevExtensions) { bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); + if (!DeviceSupports || InhibitNativeImpl) { Programs.push_back( loadDeviceLibFallback(Context, Ext, Devices, CachedLibPrograms)); - break; + FallbackIsLoaded = true; } } } + return Programs; } ProgramManager::ProgramPtr @@ -486,7 +493,7 @@ ProgramManager::build(ProgramPtr Program, RT::PiContext Context, std::vector LinkPrograms; if (LinkDeviceLibs) { - getDeviceLibPrograms(Context, Devices, CachedLibPrograms, LinkPrograms); + LinkPrograms = getDeviceLibPrograms(Context, Devices, CachedLibPrograms); } if (LinkPrograms.empty()) { From 27bcf97e066324c1abb41be3a1af1e58c553b233 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 15:06:59 +0300 Subject: [PATCH 24/32] Use an enum instead of std::string for extension id Signed-off-by: Andrew Savonichev --- sycl/include/CL/sycl/detail/context_impl.hpp | 11 ++-- .../program_manager/program_manager.hpp | 6 ++- .../program_manager/program_manager.cpp | 53 ++++++++++++------- 3 files changed, 44 insertions(+), 26 deletions(-) diff --git a/sycl/include/CL/sycl/detail/context_impl.hpp b/sycl/include/CL/sycl/detail/context_impl.hpp index 983b3abffe681..4e316671079b5 100644 --- a/sycl/include/CL/sycl/detail/context_impl.hpp +++ b/sycl/include/CL/sycl/detail/context_impl.hpp @@ -11,6 +11,7 @@ #include #include #include +#include #include #include #include @@ -130,17 +131,17 @@ class context_impl { /// 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 name: + /// corresponding extension: /// - /// "cl_intel_devicelib_assert" -> # - /// "cl_intel_devicelib_complex" -> # + /// cl_intel_devicelib_assert -> # + /// cl_intel_devicelib_complex -> # /// etc. /// /// See `doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst' for /// more details. /// /// @returns a map with device library programs. - std::map &getCachedLibPrograms() { + std::map &getCachedLibPrograms() { return MCachedLibPrograms; } @@ -154,7 +155,7 @@ class context_impl { std::map MCachedPrograms; std::map> MCachedKernels; std::shared_ptr MUSMDispatch; - std::map MCachedLibPrograms; + std::map MCachedLibPrograms; }; } // namespace detail diff --git a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp b/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp index ef6b0a25aa2e3..1774c90cbe195 100644 --- a/sycl/include/CL/sycl/detail/program_manager/program_manager.hpp +++ b/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 { @@ -75,7 +79,7 @@ class ProgramManager { ProgramPtr build(ProgramPtr Program, RT::PiContext Context, const string_class &Options, const std::vector &Devices, - std::map &CachedLibPrograms, + std::map &CachedLibPrograms, bool LinkDeviceLibs = false); /// Provides a new kernel set id for grouping kernel names together KernelSetId getNextKernelSetId() const; diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index d9ec59df49201..4a9eeb5d3af2d 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -287,20 +287,31 @@ static std::string getDeviceExtensions(const RT::PiDevice &Dev) { return DevExt; } -static RT::PiProgram -loadDeviceLibFallback(const RT::PiContext &Context, - const std::string &Extension, - const std::vector &Devices, - std::map &CachedLibPrograms) { - - const char *LibFileName = nullptr; - if (Extension == "cl_intel_devicelib_assert") { - LibFileName = "libsycl-fallback-cassert.spv"; - } else { - throw compile_program_error(std::string("Unknown device library: ") + - Extension); +static const char* getDeviceLibFilename(DeviceLibExt Extension) { + switch (Extension) { + case cl_intel_devicelib_assert: + return "libsycl-fallback-cassert.spv"; + } + throw compile_program_error("Unhandled (new?) device library extension"); +} + +const char* getDeviceLibExtensionStr(DeviceLibExt Extension) { + switch (Extension) { + case cl_intel_devicelib_assert: + return "cl_intel_devicelib_assert"; } - std::map::iterator LibProgIt; + throw compile_program_error("Unhandled (new?) device library extension"); +} + +static RT::PiProgram +loadDeviceLibFallback( + const RT::PiContext &Context, + DeviceLibExt Extension, + const std::vector &Devices, + std::map &CachedLibPrograms) { + + const char *LibFileName = getDeviceLibFilename(Extension); + std::map::iterator LibProgIt; bool NotExists = false; std::tie(LibProgIt, NotExists) = CachedLibPrograms.insert({Extension, nullptr}); @@ -427,15 +438,15 @@ DeviceImage &ProgramManager::getDeviceImage(OSModuleHandle M, KernelSetId KSId, static std::vector getDeviceLibPrograms( const RT::PiContext Context, const std::vector &Devices, - std::map &CachedLibPrograms) { + std::map &CachedLibPrograms) { std::vector Programs; // TODO: SYCL compiler should generate a list of required extensions for a // particular program in order to allow us do a more fine-grained check here. // Require *all* possible devicelib extensions for now. - std::pair RequiredDeviceLibExt[] = { - {"cl_intel_devicelib_assert", false} + std::pair RequiredDeviceLibExt[] = { + {cl_intel_devicelib_assert, false} }; // Load a fallback library for an extension if at least one device does not @@ -443,19 +454,21 @@ static std::vector getDeviceLibPrograms( for (RT::PiDevice Dev : Devices) { std::string DevExtList = getDeviceExtensions(Dev); for (auto &Pair : RequiredDeviceLibExt) { - const char *Ext = Pair.first; + DeviceLibExt Ext = Pair.first; bool &FallbackIsLoaded = Pair.second; + const char* ExtStr = getDeviceLibExtensionStr(Ext); + if (FallbackIsLoaded) { continue; } bool InhibitNativeImpl = false; if (const char *Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { - InhibitNativeImpl = strstr(Env, Ext) != nullptr; + InhibitNativeImpl = strstr(Env, ExtStr) != nullptr; } - bool DeviceSupports = DevExtList.npos != DevExtList.find(Ext); + bool DeviceSupports = DevExtList.npos != DevExtList.find(ExtStr); if (!DeviceSupports || InhibitNativeImpl) { Programs.push_back( @@ -471,7 +484,7 @@ ProgramManager::ProgramPtr ProgramManager::build(ProgramPtr Program, RT::PiContext Context, const string_class &Options, const std::vector &Devices, - std::map &CachedLibPrograms, + std::map &CachedLibPrograms, bool LinkDeviceLibs) { if (DbgProgMgr > 0) { From 45b9be8d0f8343495dc93ef535337c8f1d59fd6b Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 15:14:52 +0300 Subject: [PATCH 25/32] Fix typo Signed-off-by: Andrew Savonichev --- sycl/test/devicelib/assert.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp index aa6b9702b1f1e..464a0760b7513 100644 --- a/sycl/test/devicelib/assert.cpp +++ b/sycl/test/devicelib/assert.cpp @@ -4,7 +4,7 @@ // (see the other RUN lines below; it is a bit complicated) // // assert() call in device code guarantees nothing: on some devices it behaves -// in the usual way and terminate the program. On other devices it can print an +// in the usual way and terminates the program. On other devices it can print an // error message and *continue* execution. Less capable devices can even ignore // an assert! // From 9f1aeccc79d3a2b2a225fd0b18efc0722839e456 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 15:36:45 +0300 Subject: [PATCH 26/32] Use detail::get_device_info to get device extensions Signed-off-by: Andrew Savonichev --- .../detail/program_manager/program_manager.cpp | 16 ++-------------- 1 file changed, 2 insertions(+), 14 deletions(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 321963c285e04..3dd180b35160e 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -274,19 +274,6 @@ static bool loadDeviceLib(const RT::PiContext &Context, const char *Name, return Prog != nullptr; } -static std::string getDeviceExtensions(const RT::PiDevice &Dev) { - std::string DevExt; - size_t DevExtSize = 0; - PI_CALL(piDeviceGetInfo)(Dev, PI_DEVICE_INFO_EXTENSIONS, - /*param_value_size=*/0, - /*param_value=*/nullptr, &DevExtSize); - DevExt.resize(DevExtSize); - PI_CALL(piDeviceGetInfo)(Dev, PI_DEVICE_INFO_EXTENSIONS, DevExt.size(), - &DevExt[0], - /*param_value_size_ret=*/nullptr); - return DevExt; -} - static const char* getDeviceLibFilename(DeviceLibExt Extension) { switch (Extension) { case cl_intel_devicelib_assert: @@ -452,7 +439,8 @@ static std::vector getDeviceLibPrograms( // Load a fallback library for an extension if at least one device does not // support it. for (RT::PiDevice Dev : Devices) { - std::string DevExtList = getDeviceExtensions(Dev); + std::string DevExtList = + get_device_info::get(Dev); for (auto &Pair : RequiredDeviceLibExt) { DeviceLibExt Ext = Pair.first; bool &FallbackIsLoaded = Pair.second; From 8e826ef503c262b48727b1e75e3fa067b6afe914 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 17:02:21 +0300 Subject: [PATCH 27/32] Fix typo Signed-off-by: Andrew Savonichev --- sycl/source/detail/program_manager/program_manager.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 3dd180b35160e..675a258024f8b 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -187,7 +187,7 @@ ProgramManager::getBuiltPIProgram(OSModuleHandle M, const context &Context, context_impl *ContextImpl = getRawSyclObjImpl(Context); RT::PiContext PiContext = ContextImpl->getHandleRef(); - const std::vector Devices = ContextImpl->getDevices(); + const std::vector &Devices = ContextImpl->getDevices(); std::vector PiDevices(Devices.size()); std::transform( Devices.begin(), Devices.end(), PiDevices.begin(), From cf819f18d2d766784016a91a916c74c2f327571f Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 17:09:12 +0300 Subject: [PATCH 28/32] Use uint64_t instead of size_t for __devicelib_assert_fail Signed-off-by: Andrew Savonichev --- sycl/source/detail/devicelib/fallback-cassert.cpp | 13 +++++++------ sycl/source/detail/devicelib/wrapper.h | 9 +++++---- 2 files changed, 12 insertions(+), 10 deletions(-) diff --git a/sycl/source/detail/devicelib/fallback-cassert.cpp b/sycl/source/detail/devicelib/fallback-cassert.cpp index 204741086fdb2..0c2d8ba5fa96d 100644 --- a/sycl/source/detail/devicelib/fallback-cassert.cpp +++ b/sycl/source/detail/devicelib/fallback-cassert.cpp @@ -17,10 +17,11 @@ static const __attribute__((opencl_constant)) char assert_fmt[] = "Assertion `%s` failed.\n"; SYCL_EXTERNAL -extern "C" void __devicelib_assert_fail(const char *expr, const char *file, - int32_t line, const char *func, - size_t gid0, size_t gid1, size_t gid2, - size_t lid0, size_t lid1, size_t lid2) { +extern "C" void __devicelib_assert_fail( + const char *expr, const char *file, + int32_t line, const char *func, + uint64_t gid0, uint64_t gid1, uint64_t gid2, + uint64_t lid0, uint64_t lid1, uint64_t lid2) { // intX_t types are used instead of `int' and `long' because the format string // is defined in terms of *device* types (OpenCL types): %d matches a 32 bit // integer, %lu matches a 64 bit unsigned integer. Host `int' and @@ -29,8 +30,8 @@ extern "C" void __devicelib_assert_fail(const char *expr, const char *file, assert_fmt, file, (int32_t)line, (func) ? func : "", - (uint64_t)gid0, (uint64_t)gid1, (uint64_t)gid2, - (uint64_t)lid0, (uint64_t)lid1, (uint64_t)lid2, + gid0, gid1, gid2, + lid0, lid1, lid2, expr); // FIXME: call SPIR-V unreachable instead diff --git a/sycl/source/detail/devicelib/wrapper.h b/sycl/source/detail/devicelib/wrapper.h index 6148f0436498b..b541cb9880c6c 100644 --- a/sycl/source/detail/devicelib/wrapper.h +++ b/sycl/source/detail/devicelib/wrapper.h @@ -13,9 +13,10 @@ #include SYCL_EXTERNAL -extern "C" void __devicelib_assert_fail(const char *expr, const char *file, - int32_t line, const char *func, - size_t gid0, size_t gid1, size_t gid2, - size_t lid0, size_t lid1, size_t lid2); +extern "C" void __devicelib_assert_fail( + const char *expr, const char *file, + int32_t line, const char *func, + uint64_t gid0, uint64_t gid1, uint64_t gid2, + uint64_t lid0, uint64_t lid1, uint64_t lid2); #endif // __SYCL_WRAPPERS_H__ From 8b41e4a1c8a4eb04cd0747d1019d34ac35e0822e Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 18:04:33 +0300 Subject: [PATCH 29/32] Code review comments Signed-off-by: Andrew Savonichev --- sycl/doc/SYCLEnvironmentVariables.md | 1 + sycl/source/detail/program_manager/program_manager.cpp | 8 ++++---- sycl/test/devicelib/assert.cpp | 2 +- 3 files changed, 6 insertions(+), 5 deletions(-) diff --git a/sycl/doc/SYCLEnvironmentVariables.md b/sycl/doc/SYCLEnvironmentVariables.md index 53198ec2eb525..fb36ce4329bdb 100644 --- a/sycl/doc/SYCLEnvironmentVariables.md +++ b/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.` diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index 675a258024f8b..4b890e68a2c0c 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -282,7 +282,7 @@ static const char* getDeviceLibFilename(DeviceLibExt Extension) { throw compile_program_error("Unhandled (new?) device library extension"); } -const char* getDeviceLibExtensionStr(DeviceLibExt Extension) { +static const char* getDeviceLibExtensionStr(DeviceLibExt Extension) { switch (Extension) { case cl_intel_devicelib_assert: return "cl_intel_devicelib_assert"; @@ -433,7 +433,7 @@ static std::vector getDeviceLibPrograms( // particular program in order to allow us do a more fine-grained check here. // Require *all* possible devicelib extensions for now. std::pair RequiredDeviceLibExt[] = { - {cl_intel_devicelib_assert, false} + {cl_intel_devicelib_assert, /* is fallback loaded? */ false} }; // Load a fallback library for an extension if at least one device does not @@ -445,12 +445,12 @@ static std::vector getDeviceLibPrograms( DeviceLibExt Ext = Pair.first; bool &FallbackIsLoaded = Pair.second; - const char* ExtStr = getDeviceLibExtensionStr(Ext); - if (FallbackIsLoaded) { continue; } + const char* ExtStr = getDeviceLibExtensionStr(Ext); + bool InhibitNativeImpl = false; if (const char *Env = getenv("SYCL_DEVICELIB_INHIBIT_NATIVE")) { InhibitNativeImpl = strstr(Env, ExtStr) != nullptr; diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp index 464a0760b7513..97bc1138882b5 100644 --- a/sycl/test/devicelib/assert.cpp +++ b/sycl/test/devicelib/assert.cpp @@ -56,7 +56,7 @@ // the message can still be buffered by stdio. We turn the bufferization // off explicitly. // -// SYCL_DEVICELIB_LINK_FALLBACK=1 environment variable is used to force a mode +// SYCL_DEVICELIB_INHIBIT_NATIVE=1 environment variable is used to force a mode // in SYCL Runtime, so it doesn't look into a device extensions list and always // link the fallback library. // From 423a054000910910577931a79c3fcbe321782f32 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Wed, 18 Dec 2019 18:05:09 +0300 Subject: [PATCH 30/32] Workaround bug in IGC Signed-off-by: Andrew Savonichev --- sycl/source/detail/devicelib/fallback-cassert.cpp | 4 +++- sycl/test/devicelib/assert-windows.cpp | 2 +- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/source/detail/devicelib/fallback-cassert.cpp b/sycl/source/detail/devicelib/fallback-cassert.cpp index 0c2d8ba5fa96d..0d1fc161cd10f 100644 --- a/sycl/source/detail/devicelib/fallback-cassert.cpp +++ b/sycl/source/detail/devicelib/fallback-cassert.cpp @@ -29,7 +29,9 @@ extern "C" void __devicelib_assert_fail( __spirv_ocl_printf( assert_fmt, file, (int32_t)line, - (func) ? func : "", + // WORKAROUND: IGC does not handle this well + // (func) ? func : "", + func, gid0, gid1, gid2, lid0, lid1, lid2, expr); diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp index fe2700a1b5382..74837effa52f7 100644 --- a/sycl/test/devicelib/assert-windows.cpp +++ b/sycl/test/devicelib/assert-windows.cpp @@ -16,7 +16,7 @@ // // RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_INHIBIT_NATIVE=1 SYCL_DEVICE_TYPE=CPU %t.out >%t.stdout.fallback 2>%t.stderr.fallback // RUN: FileCheck %s --check-prefix=CHECK-MESSAGE --input-file %t.stdout.fallback -// CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: : local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. +// CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: (null): local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. // // RUN: FileCheck %s --input-file %t.stdout.fallback --check-prefix=CHECK-FALLBACK // CHECK-FALLBACK: ---> piProgramLink From 42d52b8cb94ef9c7d68d5c8667b269bbe264fd81 Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 19 Dec 2019 19:06:23 +0300 Subject: [PATCH 31/32] Re-order local and global id in the assert format string Signed-off-by: Andrew Savonichev --- .../extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst | 2 +- sycl/source/detail/devicelib/fallback-cassert.cpp | 2 +- sycl/test/devicelib/assert-windows.cpp | 2 +- sycl/test/devicelib/assert.cpp | 2 +- 4 files changed, 4 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst index c368ba5d60bd0..ef896f2f55adc 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/DeviceLibExtensions.rst @@ -31,4 +31,4 @@ Arguments: Example of a message: .. code: - foo.cpp:42: void foo(int): local id: [0,0,0], global id: [0,0,0] Assertion `buf[wiID] == 0 && "Invalid value"` failed. + foo.cpp:42: void foo(int): global id: [0,0,0], local id: [0,0,0] Assertion `buf[wiID] == 0 && "Invalid value"` failed. diff --git a/sycl/source/detail/devicelib/fallback-cassert.cpp b/sycl/source/detail/devicelib/fallback-cassert.cpp index 0d1fc161cd10f..d8e9ddd0a079d 100644 --- a/sycl/source/detail/devicelib/fallback-cassert.cpp +++ b/sycl/source/detail/devicelib/fallback-cassert.cpp @@ -13,7 +13,7 @@ int __spirv_ocl_printf(const __attribute__((opencl_constant)) char *fmt, ...) __attribute__((format(printf, 1, 2))); static const __attribute__((opencl_constant)) char assert_fmt[] = - "%s:%d: %s: local id: [%lu,%lu,%lu], global id: [%lu,%lu,%lu] " + "%s:%d: %s: global id: [%lu,%lu,%lu], local id: [%lu,%lu,%lu] " "Assertion `%s` failed.\n"; SYCL_EXTERNAL diff --git a/sycl/test/devicelib/assert-windows.cpp b/sycl/test/devicelib/assert-windows.cpp index 74837effa52f7..c65a8bd6a3f45 100644 --- a/sycl/test/devicelib/assert-windows.cpp +++ b/sycl/test/devicelib/assert-windows.cpp @@ -16,7 +16,7 @@ // // RUN: env SYCL_PI_TRACE=1 SYCL_DEVICELIB_INHIBIT_NATIVE=1 SYCL_DEVICE_TYPE=CPU %t.out >%t.stdout.fallback 2>%t.stderr.fallback // RUN: FileCheck %s --check-prefix=CHECK-MESSAGE --input-file %t.stdout.fallback -// CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: (null): local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. +// CHECK-MESSAGE: {{.*}}assert-windows.cpp:{{[0-9]+}}: (null): global id: [{{[0-3]}},0,0], local id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. // // RUN: FileCheck %s --input-file %t.stdout.fallback --check-prefix=CHECK-FALLBACK // CHECK-FALLBACK: ---> piProgramLink diff --git a/sycl/test/devicelib/assert.cpp b/sycl/test/devicelib/assert.cpp index 97bc1138882b5..d0f18fe8cb544 100644 --- a/sycl/test/devicelib/assert.cpp +++ b/sycl/test/devicelib/assert.cpp @@ -87,7 +87,7 @@ // native modes (fallback prints to stdout, while native prints to stderr; we // already handled this difference in the RUN lines): // -// CHECK-MESSAGE: {{.*}}assert.cpp:{{[0-9]+}}: auto simple_vadd(const std::array &, const std::array &, std::array &)::(anonymous class)::operator()(cl::sycl::handler &)::(anonymous class)::operator()(cl::sycl::id<1>) const: local id: [{{[0-3]}},0,0], global id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. +// CHECK-MESSAGE: {{.*}}assert.cpp:{{[0-9]+}}: auto simple_vadd(const std::array &, const std::array &, std::array &)::(anonymous class)::operator()(cl::sycl::handler &)::(anonymous class)::operator()(cl::sycl::id<1>) const: global id: [{{[0-3]}},0,0], local id: [{{[0-3]}},0,0] Assertion `accessorC[wiID] == 0 && "Invalid value"` failed. // // Note that the work-item that hits the assert first may vary, since the order // of execution is undefined. We catch only the first one (whatever id it is). From 2f613420d772b378e9a3c620a286047536b1b50b Mon Sep 17 00:00:00 2001 From: Andrew Savonichev Date: Thu, 19 Dec 2019 19:10:06 +0300 Subject: [PATCH 32/32] Code review comments Signed-off-by: Andrew Savonichev --- sycl/doc/SYCLEnvironmentVariables.md | 2 +- .../C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst | 4 ++-- 2 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/doc/SYCLEnvironmentVariables.md b/sycl/doc/SYCLEnvironmentVariables.md index fb36ce4329bdb..700b0fac6fc6b 100644 --- a/sycl/doc/SYCLEnvironmentVariables.md +++ b/sycl/doc/SYCLEnvironmentVariables.md @@ -18,7 +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. | +| 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.` diff --git a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst index 42e3c22f52664..fff1d75cf735b 100644 --- a/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst +++ b/sycl/doc/extensions/C-CXX-StandardLibrary/C-CXX-StandardLibrary.rst @@ -5,7 +5,7 @@ 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 - or ), and the corresponding header has to be + or ), and the corresponding header has to be explicitly included in user code. Implementation requires a special device library to be linked with a @@ -23,7 +23,7 @@ or, in case of Windows: 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) + - assert macro (from or ) NOTE: only the GNU glibc and Microsoft C libraries are currently supported.