From e7fbf1aa0544b953e79aa99c3f3c3f1866a986c6 Mon Sep 17 00:00:00 2001 From: Sergey V Maslov Date: Tue, 19 May 2020 19:59:31 +0300 Subject: [PATCH 1/7] Commit includes: * New backend which implements interoperability for Level Zero. L0 interop handler getters and "make" functions for the following objects are implemented: platform, device, queue, program, accessor. * getNative() interoperability for the platform and program. Changes to ABI are non-breaking (new symbols are exported). Changed version of the sycl library accordingly. Author: Sergey V Maslov Signed-off-by: Artur Gainullin --- sycl/CMakeLists.txt | 2 +- sycl/include/CL/sycl/backend/Intel_level0.hpp | 81 ++++++++++++++++ sycl/include/CL/sycl/detail/pi.h | 9 +- sycl/include/CL/sycl/detail/pi.hpp | 5 +- sycl/include/CL/sycl/platform.hpp | 2 +- sycl/include/CL/sycl/program.hpp | 10 ++ sycl/plugins/cuda/pi_cuda.cpp | 6 ++ sycl/plugins/level_zero/pi_level0.cpp | 45 +++++++-- sycl/plugins/opencl/pi_opencl.cpp | 7 +- sycl/source/CMakeLists.txt | 1 + sycl/source/backend/Intel_level0.cpp | 96 +++++++++++++++++++ sycl/source/backend/opencl.cpp | 6 +- sycl/source/detail/device_impl.cpp | 4 +- sycl/source/detail/platform_impl.cpp | 1 + sycl/source/detail/platform_impl.hpp | 1 - sycl/source/detail/program_impl.cpp | 11 ++- sycl/source/detail/program_impl.hpp | 3 + sycl/source/program.cpp | 3 + sycl/source/queue.cpp | 2 +- sycl/test/abi/sycl_symbols_linux.dump | 6 ++ 20 files changed, 276 insertions(+), 25 deletions(-) create mode 100644 sycl/include/CL/sycl/backend/Intel_level0.hpp create mode 100644 sycl/source/backend/Intel_level0.cpp diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt index 97481dfb30f21..5c1bbdd49b5f7 100644 --- a/sycl/CMakeLists.txt +++ b/sycl/CMakeLists.txt @@ -9,7 +9,7 @@ option(SYCL_ENABLE_WERROR "Treat all warnings as errors in SYCL project" OFF) option(SYCL_ADD_DEV_VERSION_POSTFIX "Adds -V postfix to version string" ON) set(SYCL_MAJOR_VERSION 1) -set(SYCL_MINOR_VERSION 0) +set(SYCL_MINOR_VERSION 1) set(SYCL_PATCH_VERSION 0) set(SYCL_DEV_ABI_VERSION 1) if (SYCL_ADD_DEV_VERSION_POSTFIX) diff --git a/sycl/include/CL/sycl/backend/Intel_level0.hpp b/sycl/include/CL/sycl/backend/Intel_level0.hpp new file mode 100644 index 0000000000000..0914c72680226 --- /dev/null +++ b/sycl/include/CL/sycl/backend/Intel_level0.hpp @@ -0,0 +1,81 @@ +//==------- Intel_level0.hpp - SYCL Level-Zero backend ---------------------==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#pragma once + +//#include This should be included from user code +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { + +template <> struct interop { + using type = ze_driver_handle_t; +}; + +template <> struct interop { + using type = ze_device_handle_t; +}; + +template <> struct interop { + using type = ze_command_queue_handle_t; +}; + +template <> struct interop { + using type = ze_module_handle_t; +}; + +template +struct interop> { + using type = char *; +}; + +namespace level0 { + +// Implementation of varios "make" functions resides in libsycl.so +platform make_platform(pi_native_handle NativeHandle); +device make_device(const platform &Platform, pi_native_handle NativeHandle); +program make_program(const context &Context, pi_native_handle NativeHandle); +queue make_queue(const context &Context, pi_native_handle InteropHandle); + +// Construction of SYCL platform. +template ::value>::type * = nullptr> +T make(typename interop::type Interop) { + return make_platform(reinterpret_cast(Interop)); +} + +// Construction of SYCL device. +template ::value>::type * = nullptr> +T make(const platform &Platform, + typename interop::type Interop) { + return make_device(Platform, reinterpret_cast(Interop)); +} + +// Construction of SYCL program. +template ::value>::type * = nullptr> +T make(const context &Context, + typename interop::type Interop) { + return make_program(Context, reinterpret_cast(Interop)); +} + +// Construction of SYCL queue. +template ::value>::type * = nullptr> +T make(const context &Context, + typename interop::type Interop) { + return make_queue(Context, reinterpret_cast(Interop)); +} + +} // namespace level0 +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/include/CL/sycl/detail/pi.h b/sycl/include/CL/sycl/detail/pi.h index 6718261c485b8..bbc9229e23637 100644 --- a/sycl/include/CL/sycl/detail/pi.h +++ b/sycl/include/CL/sycl/detail/pi.h @@ -850,9 +850,10 @@ piextDeviceGetNativeHandle(pi_device device, pi_native_handle *nativeHandle); /// NOTE: The created PI object takes ownership of the native handle. /// /// \param nativeHandle is the native handle to create PI device from. +/// \param platform is the platform of the device. /// \param device is the PI device created from the native handle. __SYCL_EXPORT pi_result piextDeviceCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_device *device); + pi_native_handle nativeHandle, pi_platform platform, pi_device *device); /// Selects the most appropriate device binary based on runtime information /// and the IR characteristics. @@ -944,9 +945,10 @@ piextQueueGetNativeHandle(pi_queue queue, pi_native_handle *nativeHandle); /// NOTE: The created PI object takes ownership of the native handle. /// /// \param nativeHandle is the native handle to create PI queue from. +/// \param context is the PI context of the queue. /// \param queue is the PI queue created from the native handle. __SYCL_EXPORT pi_result piextQueueCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_queue *queue); + pi_native_handle nativeHandle, pi_context context, pi_queue *queue); // // Memory @@ -1066,9 +1068,10 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle); /// NOTE: The created PI object takes ownership of the native handle. /// /// \param nativeHandle is the native handle to create PI program from. +/// \param context is the PI context of the program. /// \param program is the PI program created from the native handle. __SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_program *program); + pi_native_handle nativeHandle, pi_context context, pi_program *program); // // Kernel diff --git a/sycl/include/CL/sycl/detail/pi.hpp b/sycl/include/CL/sycl/detail/pi.hpp index 5ff6e6312df36..654320722a5fb 100644 --- a/sycl/include/CL/sycl/detail/pi.hpp +++ b/sycl/include/CL/sycl/detail/pi.hpp @@ -334,14 +334,15 @@ template inline To cast(From value) { // These conversions should use PI interop API. template <> inline pi::PiProgram cast(cl_program) { - RT::assertion(false, "pi::cast -> use piextProgramFromNative"); + RT::assertion(false, "pi::cast -> use piextCreateProgramWithNativeHandle"); return {}; } template <> inline pi::PiDevice cast(cl_device_id) { - RT::assertion(false, "pi::cast -> use piextDeviceFromNative"); + RT::assertion(false, "pi::cast -> use piextCreateDeviceWithNativeHandle"); return {}; } + } // namespace pi } // namespace detail diff --git a/sycl/include/CL/sycl/platform.hpp b/sycl/include/CL/sycl/platform.hpp index 6c9d98fa8af1e..172c2ab289a79 100644 --- a/sycl/include/CL/sycl/platform.hpp +++ b/sycl/include/CL/sycl/platform.hpp @@ -107,7 +107,7 @@ class __SYCL_EXPORT platform { /// \return a native handle, the type of which defined by the backend. template auto get_native() const -> typename interop::type { - return detail::pi::cast::type>( + return reinterpret_cast::type>( getNative()); } diff --git a/sycl/include/CL/sycl/program.hpp b/sycl/include/CL/sycl/program.hpp index c08e2b454cd5a..c6dbebf3f45bf 100644 --- a/sycl/include/CL/sycl/program.hpp +++ b/sycl/include/CL/sycl/program.hpp @@ -322,7 +322,17 @@ class __SYCL_EXPORT program { #endif // __SYCL_DEVICE_ONLY__ } + /// Gets the native handle of the SYCL platform. + /// + /// \return a native handle, the type of which defined by the backend. + template + auto get_native() const -> typename interop::type { + return reinterpret_cast::type>( + getNative()); + } + private: + pi_native_handle getNative() const; program(shared_ptr_class impl); /// Template-free version of get_kernel. diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index c839ab88707df..4ee82c54579d4 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -1308,10 +1308,12 @@ pi_result cuda_piextDeviceGetNativeHandle(pi_device device, /// NOTE: The created PI object takes ownership of the native handle. /// /// \param[in] nativeHandle The native handle to create PI device object from. +/// \param[in] platform is the PI platform of the device. /// \param[out] device Set to the PI device object created from native handle. /// /// \return TBD pi_result cuda_piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_platform platform, pi_device *device) { cl::sycl::detail::pi::die( "Creation of PI device from native handle not implemented"); @@ -1845,10 +1847,12 @@ pi_result cuda_piextQueueGetNativeHandle(pi_queue queue, /// NOTE: The created PI object takes ownership of the native handle. /// /// \param[in] nativeHandle The native handle to create PI queue object from. +/// \param[in] context is the PI context of the queue. /// \param[out] queue Set to the PI queue object created from native handle. /// /// \return TBD pi_result cuda_piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, pi_queue *queue) { cl::sycl::detail::pi::die( "Creation of PI queue from native handle not implemented"); @@ -2459,10 +2463,12 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program, /// NOTE: The created PI object takes ownership of the native handle. /// /// \param[in] nativeHandle The native handle to create PI program object from. +/// \param[in] context The PI context of the program. /// \param[out] program Set to the PI program object created from native handle. /// /// \return TBD pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, pi_program *program) { cl::sycl::detail::pi::die( "Creation of PI program from native handle not implemented"); diff --git a/sycl/plugins/level_zero/pi_level0.cpp b/sycl/plugins/level_zero/pi_level0.cpp index dc4ccfd988ca9..cedd157ab6907 100644 --- a/sycl/plugins/level_zero/pi_level0.cpp +++ b/sycl/plugins/level_zero/pi_level0.cpp @@ -1181,10 +1181,16 @@ pi_result piextDeviceGetNativeHandle(pi_device Device, } pi_result piextDeviceCreateWithNativeHandle(pi_native_handle NativeHandle, + pi_platform Platform, pi_device *Device) { + assert(NativeHandle); + assert(Device); + assert(Platform); + // Create PI device from the given L0 device handle. - die("piextDeviceCreateWithNativeHandle: not supported"); - return PI_SUCCESS; + auto ZeDevice = pi_cast(NativeHandle); + *Device = new _pi_device(ZeDevice, Platform); + return (*Device)->initialize(); } pi_result piContextCreate(const pi_context_properties *Properties, @@ -1367,15 +1373,27 @@ pi_result piQueueFinish(pi_queue Queue) { return PI_SUCCESS; } + pi_result piextQueueGetNativeHandle(pi_queue Queue, pi_native_handle *NativeHandle) { - die("piextQueueGetNativeHandle: not supported"); + assert(Queue); + assert(NativeHandle); + + auto ZeQueue = pi_cast(NativeHandle); + // Extract the L0 queue handle from the given PI queue + *ZeQueue = Queue->ZeCommandQueue; return PI_SUCCESS; } pi_result piextQueueCreateWithNativeHandle(pi_native_handle NativeHandle, + pi_context Context, pi_queue *Queue) { - die("piextQueueCreateWithNativeHandle: not supported"); + assert(NativeHandle); + assert(Context); + assert(Queue); + + auto ZeQueue = pi_cast(NativeHandle); + *Queue = new _pi_queue(ZeQueue, Context); return PI_SUCCESS; } @@ -1869,13 +1887,28 @@ pi_result piProgramRelease(pi_program Program) { pi_result piextProgramGetNativeHandle(pi_program Program, pi_native_handle *NativeHandle) { - die("piextProgramGetNativeHandle: not supported"); + assert(Program); + assert(NativeHandle); + + auto ZeModule = pi_cast(NativeHandle); + // Extract the L0 module handle from the given PI program + *ZeModule = Program->ZeModule; return PI_SUCCESS; } pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle, + pi_context Context, pi_program *Program) { - die("piextProgramCreateWithNativeHandle: not supported"); + assert(NativeHandle); + assert(Context); + assert(Program); + + auto ZeModule = pi_cast(NativeHandle); + assert(*ZeModule); + // Create PI program from the given L0 module handle + auto ZePIProgram = new _pi_program(*ZeModule, Context); + + *Program = pi_cast(ZePIProgram); return PI_SUCCESS; } diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index dd51a6501d26a..718132f9a1180 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -160,7 +160,6 @@ static pi_result USMSetIndirectAccess(pi_kernel kernel) { extern "C" { -// Example of a PI interface that does not map exactly to an OpenCL one. pi_result piPlatformsGet(pi_uint32 num_entries, pi_platform *platforms, pi_uint32 *num_platforms) { cl_int result = clGetPlatformIDs(cast(num_entries), @@ -184,7 +183,6 @@ pi_result piextPlatformCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } -// Example of a PI interface that does not map exactly to an OpenCL one. pi_result piDevicesGet(pi_platform platform, pi_device_type device_type, pi_uint32 num_entries, pi_device *devices, pi_uint32 *num_devices) { @@ -274,7 +272,7 @@ pi_result piextDeviceSelectBinary(pi_device device, pi_device_binary *images, } pi_result piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_device *piDevice) { + pi_platform, pi_device *piDevice) { assert(piDevice != nullptr); *piDevice = reinterpret_cast(nativeHandle); return PI_SUCCESS; @@ -321,7 +319,7 @@ pi_result piQueueCreate(pi_context context, pi_device device, } pi_result piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_queue *piQueue) { + pi_context, pi_queue *piQueue) { assert(piQueue != nullptr); *piQueue = reinterpret_cast(nativeHandle); return PI_SUCCESS; @@ -406,6 +404,7 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length, } pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context, pi_program *piProgram) { assert(piProgram != nullptr); *piProgram = reinterpret_cast(nativeHandle); diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 004b596e84955..c9a6977b7a472 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -95,6 +95,7 @@ endfunction(add_sycl_rt_library) set(SYCL_SOURCES "${sycl_inc_dir}/CL/sycl.hpp" "backend/opencl.cpp" + "backend/Intel_level0.cpp" "detail/accessor_impl.cpp" "detail/buffer_impl.cpp" "detail/builtins_common.cpp" diff --git a/sycl/source/backend/Intel_level0.cpp b/sycl/source/backend/Intel_level0.cpp new file mode 100644 index 0000000000000..1dc21655c4c7d --- /dev/null +++ b/sycl/source/backend/Intel_level0.cpp @@ -0,0 +1,96 @@ +//==------- Intel_level0.cpp - SYCL Level-Zero backend ---------------------==// +// +// 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 +#include +#include +#include +#include + +__SYCL_INLINE_NAMESPACE(cl) { +namespace sycl { +namespace level0 { +using namespace detail; + +// Get the L0 plugin. +static const plugin &getPlugin() { + static const plugin *L0Plugin = nullptr; + if (L0Plugin) + return *L0Plugin; + + const vector_class &Plugins = pi::initialize(); + for (const auto &Plugin : Plugins) + if (Plugin.getBackend() == backend::level0) { + L0Plugin = &Plugin; + break; + } + if (!L0Plugin) { + throw runtime_error("sycl::level0 - no Level-Zero plugin", + PI_INVALID_OPERATION); + } + return *L0Plugin; +} + +//---------------------------------------------------------------------------- +// Implementation of level0::make +__SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle) { + const auto &Plugin = getPlugin(); + // Create PI platform first. + pi::PiPlatform PiPlatform; + Plugin.call(NativeHandle, + &PiPlatform); + + // Construct the SYCL platform from PI platfrom. + return detail::createSyclObjFromImpl( + std::make_shared(PiPlatform, Plugin)); +} + +//---------------------------------------------------------------------------- +// Implementation of level0::make +__SYCL_EXPORT device make_device(const platform &Platform, + pi_native_handle NativeHandle) { + const auto &Plugin = getPlugin(); + const auto &PlatformImpl = getSyclObjImpl(Platform); + // Create PI device first. + pi::PiDevice PiDevice; + Plugin.call( + NativeHandle, PlatformImpl->getHandleRef(), &PiDevice); + // Construct the SYCL device from PI device. + return detail::createSyclObjFromImpl( + std::make_shared(PiDevice, PlatformImpl)); +} + +//---------------------------------------------------------------------------- +// Implementation of level0::make +__SYCL_EXPORT program make_program(const context &Context, + pi_native_handle NativeHandle) { + // Construct the SYCL program from native program. + // TODO: move here the code that creates PI program, and remove the + // native interop constructor. + return detail::createSyclObjFromImpl( + std::make_shared(getSyclObjImpl(Context), NativeHandle)); +} + +//---------------------------------------------------------------------------- +// Implementation of level0::make +__SYCL_EXPORT queue make_queue(const context &Context, + pi_native_handle NativeHandle) { + const auto &Plugin = getPlugin(); + const auto &ContextImpl = getSyclObjImpl(Context); + // Create PI queue first. + pi::PiQueue PiQueue; + Plugin.call( + NativeHandle, ContextImpl->getHandleRef(), &PiQueue); + // Construct the SYCL queue from PI queue. + return detail::createSyclObjFromImpl(std::make_shared( + PiQueue, ContextImpl, ContextImpl->get_async_handler())); +} + +} // namespace level0 +} // namespace sycl +} // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/backend/opencl.cpp b/sycl/source/backend/opencl.cpp index 20f0add1f85d6..bd06563c0b8b8 100644 --- a/sycl/source/backend/opencl.cpp +++ b/sycl/source/backend/opencl.cpp @@ -38,7 +38,7 @@ __SYCL_EXPORT device make_device(pi_native_handle NativeHandle) { // Create PI device first. pi::PiDevice PiDevice; Plugin.call(NativeHandle, - &PiDevice); + nullptr, &PiDevice); // Construct the SYCL device from PI device. return detail::createSyclObjFromImpl( std::make_shared(PiDevice, Plugin)); @@ -76,8 +76,8 @@ __SYCL_EXPORT queue make_queue(const context &Context, const auto &ContextImpl = getSyclObjImpl(Context); // Create PI queue first. pi::PiQueue PiQueue; - Plugin.call(NativeHandle, - &PiQueue); + Plugin.call( + NativeHandle, ContextImpl->getHandleRef(), &PiQueue); // Construct the SYCL queue from PI queue. return detail::createSyclObjFromImpl(std::make_shared( PiQueue, ContextImpl, ContextImpl->get_async_handler())); diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index f060986438c63..244fc00cda9d8 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -40,8 +40,10 @@ device_impl::device_impl(pi_native_handle InteropDeviceHandle, if (Device == nullptr) { assert(InteropDeviceHandle); // Get PI device from the raw device handle. + // NOTE: this is for OpenCL interop only (and should go away). + // With SYCL-2020 BE generalization "make" functions are used instead. Plugin.call( - InteropDeviceHandle, &MDevice); + InteropDeviceHandle, nullptr, &MDevice); InteroperabilityConstructor = true; } diff --git a/sycl/source/detail/platform_impl.cpp b/sycl/source/detail/platform_impl.cpp index ed8724e92179f..e65bf5c7ccf82 100644 --- a/sycl/source/detail/platform_impl.cpp +++ b/sycl/source/detail/platform_impl.cpp @@ -9,6 +9,7 @@ #include #include #include +#include #include #include diff --git a/sycl/source/detail/platform_impl.hpp b/sycl/source/detail/platform_impl.hpp index 8f2477d567ea1..423c1e286e709 100644 --- a/sycl/source/detail/platform_impl.hpp +++ b/sycl/source/detail/platform_impl.hpp @@ -11,7 +11,6 @@ #include #include #include -#include #include #include diff --git a/sycl/source/detail/program_impl.cpp b/sycl/source/detail/program_impl.cpp index 725cc690433b3..d5201561a7225 100644 --- a/sycl/source/detail/program_impl.cpp +++ b/sycl/source/detail/program_impl.cpp @@ -110,8 +110,8 @@ program_impl::program_impl(ContextImplPtr Context, assert(InteropProgram && "No InteropProgram/PiProgram defined with piextProgramFromNative"); // Translate the raw program handle into PI program. - Plugin.call(InteropProgram, - &MProgram); + Plugin.call( + InteropProgram, MContext->getHandleRef(), &MProgram); } else Plugin.call(Program); @@ -508,6 +508,13 @@ void program_impl::flush_spec_constants(const RTDeviceBinaryImage &Img, } } +pi_native_handle program_impl::getNative() const { + const auto &Plugin = getPlugin(); + pi_native_handle Handle; + Plugin.call(MProgram, &Handle); + return Handle; +} + } // namespace detail } // namespace sycl } // __SYCL_INLINE_NAMESPACE(cl) diff --git a/sycl/source/detail/program_impl.hpp b/sycl/source/detail/program_impl.hpp index 4205e806a6a91..eee8497ad2c1f 100644 --- a/sycl/source/detail/program_impl.hpp +++ b/sycl/source/detail/program_impl.hpp @@ -318,6 +318,9 @@ class program_impl { /// Tells whether a specialization constant has been set for this program. bool hasSetSpecConstants() const { return !SpecConstRegistry.empty(); } + /// Returns the native plugin handle. + pi_native_handle getNative() const; + private: // Deligating Constructor used in Implementation. program_impl(ContextImplPtr Context, pi_native_handle InteropProgram, diff --git a/sycl/source/program.cpp b/sycl/source/program.cpp index 6884e91cc566b..c99842038c8bc 100644 --- a/sycl/source/program.cpp +++ b/sycl/source/program.cpp @@ -35,6 +35,9 @@ program::program(const context &context, cl_program clProgram) // must retain it in order to adhere to SYCL 1.2.1 spec (Rev6, section 4.3.1.) clRetainProgram(clProgram); } + +pi_native_handle program::getNative() const { return impl->getNative(); } + program::program(std::shared_ptr impl) : impl(impl) {} cl_program program::get() const { return impl->get(); } diff --git a/sycl/source/queue.cpp b/sycl/source/queue.cpp index ab4da16014beb..69a4b64712ee5 100644 --- a/sycl/source/queue.cpp +++ b/sycl/source/queue.cpp @@ -64,7 +64,7 @@ queue::queue(const device &syclDevice, const async_handler &asyncHandler, queue::queue(cl_command_queue clQueue, const context &syclContext, const async_handler &asyncHandler) { impl = std::make_shared( - detail::pi::cast(clQueue), + reinterpret_cast(clQueue), detail::getSyclObjImpl(syclContext), asyncHandler); } diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index e975f36913e2a..0156e0ec8cc26 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3539,3 +3539,9 @@ _ZNK2cl4sycl9exception11has_contextEv _ZNK2cl4sycl9exception4whatEv __sycl_register_lib __sycl_unregister_lib +_ZN2cl4sycl6level011make_deviceERKNS0_8platformEm +_ZN2cl4sycl6level010make_queueERKNS0_7contextEm +_ZN2cl4sycl6level012make_programERKNS0_7contextEm +_ZN2cl4sycl6level013make_platformEm +_ZNK2cl4sycl8platform9getNativeEv +_ZNK2cl4sycl7program9getNativeEv From 11292a9e2f72554635c73bb54e84a3d9e71d009d Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Fri, 26 Jun 2020 10:59:43 -0700 Subject: [PATCH 2/7] Minor fixes --- sycl/include/CL/sycl/backend/Intel_level0.hpp | 2 +- sycl/plugins/level_zero/pi_level0.cpp | 3 +-- 2 files changed, 2 insertions(+), 3 deletions(-) diff --git a/sycl/include/CL/sycl/backend/Intel_level0.hpp b/sycl/include/CL/sycl/backend/Intel_level0.hpp index 0914c72680226..e0cd11bc67c14 100644 --- a/sycl/include/CL/sycl/backend/Intel_level0.hpp +++ b/sycl/include/CL/sycl/backend/Intel_level0.hpp @@ -39,7 +39,7 @@ struct interop(NativeHandle); + auto ZeQueue = pi_cast(NativeHandle); // Extract the L0 queue handle from the given PI queue *ZeQueue = Queue->ZeCommandQueue; return PI_SUCCESS; From dfe80307de74fcfc7dfd8248d8a5f7ad6a55f52c Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Fri, 26 Jun 2020 12:40:54 -0700 Subject: [PATCH 3/7] Rename files --- .../CL/sycl/backend/{Intel_level0.hpp => level_zero.hpp} | 2 +- sycl/source/CMakeLists.txt | 2 +- sycl/source/backend/{Intel_level0.cpp => level_zero.cpp} | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) rename sycl/include/CL/sycl/backend/{Intel_level0.hpp => level_zero.hpp} (97%) rename sycl/source/backend/{Intel_level0.cpp => level_zero.cpp} (98%) diff --git a/sycl/include/CL/sycl/backend/Intel_level0.hpp b/sycl/include/CL/sycl/backend/level_zero.hpp similarity index 97% rename from sycl/include/CL/sycl/backend/Intel_level0.hpp rename to sycl/include/CL/sycl/backend/level_zero.hpp index e0cd11bc67c14..3bfa1a9084f9a 100644 --- a/sycl/include/CL/sycl/backend/Intel_level0.hpp +++ b/sycl/include/CL/sycl/backend/level_zero.hpp @@ -1,4 +1,4 @@ -//==------- Intel_level0.hpp - SYCL Level-Zero backend ---------------------==// +//==--------- level_zero.hpp - SYCL Level-Zero backend ---------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index c9a6977b7a472..5c202d369e2d3 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -95,7 +95,7 @@ endfunction(add_sycl_rt_library) set(SYCL_SOURCES "${sycl_inc_dir}/CL/sycl.hpp" "backend/opencl.cpp" - "backend/Intel_level0.cpp" + "backend/level_zero.cpp" "detail/accessor_impl.cpp" "detail/buffer_impl.cpp" "detail/builtins_common.cpp" diff --git a/sycl/source/backend/Intel_level0.cpp b/sycl/source/backend/level_zero.cpp similarity index 98% rename from sycl/source/backend/Intel_level0.cpp rename to sycl/source/backend/level_zero.cpp index 1dc21655c4c7d..d0b1426610706 100644 --- a/sycl/source/backend/Intel_level0.cpp +++ b/sycl/source/backend/level_zero.cpp @@ -1,4 +1,4 @@ -//==------- Intel_level0.cpp - SYCL Level-Zero backend ---------------------==// +//==--------- level_zero.cpp - SYCL Level-Zero backend ---------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 28ba02012507399fb08288f0ecb4581581c3155d Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Tue, 30 Jun 2020 16:42:18 -0700 Subject: [PATCH 4/7] Address review comments --- sycl/include/CL/sycl/backend/cuda.hpp | 2 ++ sycl/include/CL/sycl/backend/level_zero.hpp | 2 +- 2 files changed, 3 insertions(+), 1 deletion(-) diff --git a/sycl/include/CL/sycl/backend/cuda.hpp b/sycl/include/CL/sycl/backend/cuda.hpp index d215800e673b7..c266f5b715d28 100644 --- a/sycl/include/CL/sycl/backend/cuda.hpp +++ b/sycl/include/CL/sycl/backend/cuda.hpp @@ -40,6 +40,8 @@ template <> struct interop { using type = CUstream; }; template <> struct interop { using type = CUevent; }; +template <> struct interop { using type = CUmodule; }; + template struct interop This should be included from user code #include +#include __SYCL_INLINE_NAMESPACE(cl) { namespace sycl { From c7d3fd956458fa32708ec5816f038ceafd051b9d Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Tue, 30 Jun 2020 16:52:38 -0700 Subject: [PATCH 5/7] Define CUmodule type --- sycl/include/CL/sycl/backend/cuda.hpp | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/include/CL/sycl/backend/cuda.hpp b/sycl/include/CL/sycl/backend/cuda.hpp index c266f5b715d28..5235131e62401 100644 --- a/sycl/include/CL/sycl/backend/cuda.hpp +++ b/sycl/include/CL/sycl/backend/cuda.hpp @@ -21,6 +21,7 @@ typedef int CUdevice; typedef struct CUctx_st *CUcontext; typedef struct CUstream_st *CUstream; typedef struct CUevent_st *CUevent; +typedef struct CUmod_st *CUmodule; // As defined in the CUDA 10.1 header file. This requires CUDA version > 3.2 #if defined(_WIN64) || defined(__LP64__) From 9301a53d3e15693abdce6ca302c2953d54c0b099 Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Wed, 1 Jul 2020 13:09:57 -0700 Subject: [PATCH 6/7] Fix piextProgramCreateWithNativeHandle after merge --- sycl/plugins/level_zero/pi_level0.cpp | 25 ++++++++++++++++++++----- 1 file changed, 20 insertions(+), 5 deletions(-) diff --git a/sycl/plugins/level_zero/pi_level0.cpp b/sycl/plugins/level_zero/pi_level0.cpp index bf3799d9e84d7..29fb2fa8b2840 100644 --- a/sycl/plugins/level_zero/pi_level0.cpp +++ b/sycl/plugins/level_zero/pi_level0.cpp @@ -1906,12 +1906,27 @@ pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle, assert(Context); assert(Program); - auto ZeModule = pi_cast(NativeHandle); - assert(*ZeModule); - // Create PI program from the given L0 module handle - auto ZePIProgram = new _pi_program(*ZeModule, Context); + auto ZeModule = pi_cast(NativeHandle); + + // Create PI program from the given L0 module handle. + // + // TODO: We don't have the real L0 module descriptor with + // which it was created, but that's only needed for zeModuleCreate, + // which we don't expect to be called on the interop program. + // + ze_module_desc_t ZeModuleDesc = {}; + ZeModuleDesc.version = ZE_MODULE_DESC_VERSION_CURRENT; + ZeModuleDesc.format = ZE_MODULE_FORMAT_NATIVE; + ZeModuleDesc.inputSize = 0; + ZeModuleDesc.pInputModule = nullptr; - *Program = pi_cast(ZePIProgram); + try { + *Program = new _pi_program(ZeModule, ZeModuleDesc, Context); + } catch (const std::bad_alloc &) { + return PI_OUT_OF_HOST_MEMORY; + } catch (...) { + return PI_ERROR_UNKNOWN; + } return PI_SUCCESS; } From a664a0e42081f82547e47261d87e433f19967fc9 Mon Sep 17 00:00:00 2001 From: Artur Gainullin Date: Wed, 1 Jul 2020 13:18:32 -0700 Subject: [PATCH 7/7] Reuse pi::getPlugin --- sycl/source/backend/level_zero.cpp | 25 +++---------------------- sycl/source/detail/pi.cpp | 1 + 2 files changed, 4 insertions(+), 22 deletions(-) diff --git a/sycl/source/backend/level_zero.cpp b/sycl/source/backend/level_zero.cpp index d0b1426610706..2e62223c0301e 100644 --- a/sycl/source/backend/level_zero.cpp +++ b/sycl/source/backend/level_zero.cpp @@ -17,29 +17,10 @@ namespace sycl { namespace level0 { using namespace detail; -// Get the L0 plugin. -static const plugin &getPlugin() { - static const plugin *L0Plugin = nullptr; - if (L0Plugin) - return *L0Plugin; - - const vector_class &Plugins = pi::initialize(); - for (const auto &Plugin : Plugins) - if (Plugin.getBackend() == backend::level0) { - L0Plugin = &Plugin; - break; - } - if (!L0Plugin) { - throw runtime_error("sycl::level0 - no Level-Zero plugin", - PI_INVALID_OPERATION); - } - return *L0Plugin; -} - //---------------------------------------------------------------------------- // Implementation of level0::make __SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle) { - const auto &Plugin = getPlugin(); + const auto &Plugin = pi::getPlugin(); // Create PI platform first. pi::PiPlatform PiPlatform; Plugin.call(NativeHandle, @@ -54,7 +35,7 @@ __SYCL_EXPORT platform make_platform(pi_native_handle NativeHandle) { // Implementation of level0::make __SYCL_EXPORT device make_device(const platform &Platform, pi_native_handle NativeHandle) { - const auto &Plugin = getPlugin(); + const auto &Plugin = pi::getPlugin(); const auto &PlatformImpl = getSyclObjImpl(Platform); // Create PI device first. pi::PiDevice PiDevice; @@ -80,7 +61,7 @@ __SYCL_EXPORT program make_program(const context &Context, // Implementation of level0::make __SYCL_EXPORT queue make_queue(const context &Context, pi_native_handle NativeHandle) { - const auto &Plugin = getPlugin(); + const auto &Plugin = pi::getPlugin(); const auto &ContextImpl = getSyclObjImpl(Context); // Create PI queue first. pi::PiQueue PiQueue; diff --git a/sycl/source/detail/pi.cpp b/sycl/source/detail/pi.cpp index b00844cf85a59..bfdc589b8a267 100644 --- a/sycl/source/detail/pi.cpp +++ b/sycl/source/detail/pi.cpp @@ -390,6 +390,7 @@ template const plugin &getPlugin() { } template const plugin &getPlugin(); +template const plugin &getPlugin(); // Report error and no return (keeps compiler from printing warnings). // TODO: Probably change that to throw a catchable exception,