Skip to content

[SYCL] Add Level-Zero interop with specification of ownership for kernel and kernel_bundle. #4542

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Closed
wants to merge 3 commits into from
Closed

[SYCL] Add Level-Zero interop with specification of ownership for kernel and kernel_bundle. #4542

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
3 commits
Select commit Hold shift + click to select a range
d260be3
[SYCL] Add Level-Zero interop with specification of ownership for
rbegam Sep 10, 2021
2e72d5c
adds level-zero interop with specification of ownership for kernel.
rbegam Sep 10, 2021
c71132b
clang formatted.
rbegam Sep 11, 2021
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -90,7 +90,8 @@ a SYCL object that encapsulates a corresponding Level-Zero object:
|``` make<context>(const vector_class<device> &, ze_context_handle_t, ownership = transfer);```| Constructs a SYCL context instance from a Level-Zero ```ze_context_handle_t```. The context is created against the devices passed in. There must be at least one device given and all the devices must be from the same SYCL platform and thus from the same Level-Zero driver. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.|
|``` make<queue>(const context &, ze_command_queue_handle_t, ownership = transfer);```| Constructs a SYCL queue instance from a Level-Zero ```ze_command_queue_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The queue is attached to the first device in the passed SYCL context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.|
|``` make<event>(const context &, ze_event_handle_t, ownership = transfer);```| Constructs a SYCL event instance from a Level-Zero ```ze_event_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero event should be allocated from an event pool created in the same context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.|
|``` make<program>(const context &, ze_module_handle_t);```| Constructs a SYCL program instance from a Level-Zero ```ze_module_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero module must be fully linked (i.e. not require further linking through [```zeModuleDynamicLink```](https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t)), and thus the SYCL program is created in the "linked" state.|
|``` make<kernel>(const context &, ze_kernel_handle_t, ownership = transfer);```| Constructs a SYCL kernel from a Level-Zero ```ze_kernel_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.|
|``` make<kernel_bundle>(const context &, ze_module_handle_t, ownership = transfer);```| Constructs a SYCL program instance from a Level-Zero ```ze_module_handle_t```. The context argument must be a valid SYCL context encapsulating a Level-Zero context. The Level-Zero module must be fully linked (i.e. not require further linking through [```zeModuleDynamicLink```](https://spec.oneapi.com/level-zero/latest/core/api.html?highlight=zemoduledynamiclink#_CPPv419zeModuleDynamicLink8uint32_tP18ze_module_handle_tP28ze_module_build_log_handle_t)), and thus the SYCL program is created in the "linked" state. The ```ownership``` argument specifies if the SYCL runtime should take ownership of the passed native handle. The default behavior is to transfer the ownership to the SYCL runtime. See section 4.4 for details.|
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I have several comments about these APIs:

  • It seems like make<kernel>() might also need a ze_module_handle_t parameter. There is a SYCL API kernel::get_kernel_bundle() which returns the kernel bundle that contains a kernel. How will you implement this if you don't know the underlying Level Zero module handle?

  • Both APIs should be more precise about the requirements of the context parameter. For make<kernel_bundle>(), I think it is required that the ze_module_handle_t be created from the same context. Likewise, for make<kernel>(), I think it is required that the ze_kernel_handle_t come from a module that is created from that context.

  • I think we probably need to document another requirement for make<kernel_bundle>() when ownership is transfer. Since the runtime takes ownership of the L0 handle in this case, I think the application must also promise not to have any outstanding ze_kernel_handle_t handles to the underlying ze_module_handle_t. Otherwise, the runtime won't be able to deallocate that ze_module_handle_t.

  • The prototype for make<kernel_bundle> isn't correct because kernel_bundle is a template type. Instead, I think you want to say make<kernel_bundle<bundle_state::executable>>. You can then remove the wording in the description "and thus the SYCL program is created in the "linked" state" because this is clear from the prototype.

  • The wording for make<kernel_bundle> should be updated: "Constructs a SYCL program instance". Instead say "Construct a SYCL kernel bundle instance".

I also have a general word of caution about this API. We made a decision for the last release that this document should describe the existing API even thought it does not conform to the SYCL 2020 spec, with the understanding that we would update the extension document and the API implementation to conform "soon". We should do that sooner, rather than later. All of these APIs will end up getting deprecated because none of them conform to the spec, so it would be better to migrate to a conformant API before adding too many more additions. What do you think @smaslov-intel ?

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

My recent commit #4512 makes all of the current Level-Zero interop following SYCL-2020 and deprecates prior functions. Update of https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/LevelZeroBackend/LevelZeroBackend.md is WIP. Further additions like in this PR should follow SYCL-2020 from the beginning.


NOTE: We shall consider adding other interoperability as needed, if possible.

Expand Down Expand Up @@ -197,3 +198,5 @@ struct free_memory {
|4|2021-07-06|Rehana Begam|Introduced explicit ownership for queue
|5|2021-07-25|Sergey Maslov|Introduced SYCL interop for events
|6|2021-08-30|Dmitry Vodopyanov|Updated according to SYCL 2020 reqs for extensions
|7|2021-09-09|Rehana Begam|Introduced explicit ownership for kernel_bundle
|8|2021-09-10|Rehana Begam|Introduced explicit ownership for kernel
33 changes: 31 additions & 2 deletions sycl/include/CL/sycl/backend.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -108,6 +108,12 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
backend Backend);
__SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle,
const context &TargetContext, backend Backend);
__SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle,
const context &TargetContext,
bool KeepOwnership, backend Backend);
__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
bool KeepOwnership, bundle_state State, backend Backend);
__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
bundle_state State, backend Backend);
Expand Down Expand Up @@ -199,13 +205,36 @@ make_buffer(const typename backend_traits<Backend>::template input_type<
reinterpret_cast<cl_mem>(BackendObject), TargetContext, AvailableEvent);
}

template <backend Backend>
kernel
make_kernel(const typename backend_traits<Backend>::template input_type<kernel>
&BackendObject,
const context &TargetContext, bool KeepOwnership) {
return detail::make_kernel(detail::pi::cast<pi_native_handle>(BackendObject),
TargetContext, KeepOwnership, Backend);
}

template <backend Backend>
kernel
make_kernel(const typename backend_traits<Backend>::template input_type<kernel>
&BackendObject,
const context &TargetContext) {
return detail::make_kernel(detail::pi::cast<pi_native_handle>(BackendObject),
TargetContext, Backend);
TargetContext, false, Backend);
}

template <backend Backend, bundle_state State>
typename std::enable_if<
detail::InteropFeatureSupportMap<Backend>::MakeKernelBundle == true,
kernel_bundle<State>>::type
make_kernel_bundle(const typename backend_traits<Backend>::template input_type<
kernel_bundle<State>> &BackendObject,
const context &TargetContext, bool KeepOwnership) {
std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
detail::make_kernel_bundle(
detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
KeepOwnership, State, Backend);
return detail::createSyclObjFromImpl<kernel_bundle<State>>(KBImpl);
}

template <backend Backend, bundle_state State>
Expand All @@ -218,7 +247,7 @@ make_kernel_bundle(const typename backend_traits<Backend>::template input_type<
std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
detail::make_kernel_bundle(
detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
State, Backend);
false, State, Backend);
return detail::createSyclObjFromImpl<kernel_bundle<State>>(KBImpl);
}
} // namespace sycl
Expand Down
20 changes: 14 additions & 6 deletions sycl/include/CL/sycl/detail/pi.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -40,10 +40,14 @@
// changes the API version from 3.5 to 4.6.
// 5.7 Added new context and ownership arguments to
// piextEventCreateWithNativeHandle
// 6.8 Added new ownership argument to piextProgramCreateWithNativeHandle which
// changes the API version from 5.7 to 6.8
// 7.9 Added new ownership argument to piextKernelCreateWithNativeHandle which
// changes the API version from 6.8 to 7.9
//
#include "CL/cl.h"
#define _PI_H_VERSION_MAJOR 5
#define _PI_H_VERSION_MINOR 7
#define _PI_H_VERSION_MAJOR 7
#define _PI_H_VERSION_MINOR 9

#define _PI_STRING_HELPER(a) #a
#define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b)
Expand Down Expand Up @@ -1220,8 +1224,11 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle);
/// \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.
/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
/// the native handle, if it can.
__SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle(
pi_native_handle nativeHandle, pi_context context, pi_program *program);
pi_native_handle nativeHandle, pi_context context, pi_program *program,
bool ownNativeHandle);

//
// Kernel
Expand Down Expand Up @@ -1315,12 +1322,13 @@ __SYCL_EXPORT pi_result piKernelSetExecInfo(pi_kernel kernel,
///
/// \param nativeHandle is the native handle to create PI kernel from.
/// \param context is the PI context of the kernel.
/// \param kernel is the PI kernel created from the native handle.
/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
/// the native handle, if it can.
/// \param kernel is the PI kernel created from the native handle.
///
__SYCL_EXPORT pi_result piextKernelCreateWithNativeHandle(
pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
pi_kernel *kernel);
pi_native_handle nativeHandle, pi_context context, pi_kernel *kernel,
bool ownNativeHandle);

/// Gets the native handle of a PI kernel object.
///
Expand Down
8 changes: 6 additions & 2 deletions sycl/plugins/cuda/pi_cuda.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2681,7 +2681,7 @@ pi_result cuda_piEnqueueNativeKernel(pi_queue, void (*)(void *), void *, size_t,
}

pi_result cuda_piextKernelCreateWithNativeHandle(pi_native_handle, pi_context,
bool, pi_kernel *) {
pi_kernel *, bool) {
sycl::detail::pi::die("Unsupported operation");
return PI_SUCCESS;
}
Expand Down Expand Up @@ -3161,10 +3161,14 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program,
/// \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.
/// \param[in] ownNativeHandle tells if SYCL RT should assume the ownership of
/// the native handle, if it can.
///
/// \return TBD
pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
pi_program *) {
pi_program *,
bool ownNativeHandle) {
(void)ownNativeHandle;
cl::sycl::detail::pi::die(
"Creation of PI program from native handle not implemented");
return {};
Expand Down
6 changes: 3 additions & 3 deletions sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1027,7 +1027,7 @@ pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *) {
}

pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
pi_program *) {
pi_program *, bool) {
DIE_NO_IMPLEMENTATION;
}

Expand Down Expand Up @@ -1397,8 +1397,8 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim,
return PI_SUCCESS;
}

pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, bool,
pi_kernel *) {
pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context,
pi_kernel *, bool) {
DIE_NO_IMPLEMENTATION;
}

Expand Down
20 changes: 12 additions & 8 deletions sycl/plugins/level_zero/pi_level_zero.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3235,7 +3235,7 @@ pi_result piProgramCreate(pi_context Context, const void *ILBytes,
// and piProgramCompile. Also it is only then we know the build options.

try {
*Program = new _pi_program(Context, ILBytes, Length, _pi_program::IL);
*Program = new _pi_program(Context, ILBytes, Length, _pi_program::IL, true);
} catch (const std::bad_alloc &) {
return PI_OUT_OF_HOST_MEMORY;
} catch (...) {
Expand Down Expand Up @@ -3281,7 +3281,8 @@ pi_result piProgramCreateWithBinary(
// information to distinguish the cases.

try {
*Program = new _pi_program(Context, Binary, Length, _pi_program::Native);
*Program =
new _pi_program(Context, Binary, Length, _pi_program::Native, true);
} catch (const std::bad_alloc &) {
return PI_OUT_OF_HOST_MEMORY;
} catch (...) {
Expand Down Expand Up @@ -3528,7 +3529,7 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
return res;
}
Input = new _pi_program(Input->Context, ZeModule, _pi_program::Object,
Input->HasImports);
true, Input->HasImports);
Input->HasImportsAndIsLinked = true;
}
} else {
Expand All @@ -3551,7 +3552,8 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
// the description of the failure).
if (ZeResult == ZE_RESULT_SUCCESS ||
ZeResult == ZE_RESULT_ERROR_MODULE_LINK_FAILURE) {
*RetProgram = new _pi_program(Context, std::move(Inputs), ZeBuildLog);
*RetProgram =
new _pi_program(Context, std::move(Inputs), ZeBuildLog, true);
}
if (ZeResult != ZE_RESULT_SUCCESS)
return mapError(ZeResult);
Expand Down Expand Up @@ -3783,7 +3785,8 @@ pi_result piextProgramGetNativeHandle(pi_program Program,

pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
pi_context Context,
pi_program *Program) {
pi_program *Program,
bool OwnNativeHandle) {
PI_ASSERT(Program, PI_INVALID_PROGRAM);
PI_ASSERT(NativeHandle, PI_INVALID_VALUE);
PI_ASSERT(Context, PI_INVALID_CONTEXT);
Expand All @@ -3795,7 +3798,8 @@ pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
// executable (state Object).

try {
*Program = new _pi_program(Context, ZeModule, _pi_program::Exe);
*Program =
new _pi_program(Context, ZeModule, _pi_program::Exe, OwnNativeHandle);
} catch (const std::bad_alloc &) {
return PI_OUT_OF_HOST_MEMORY;
} catch (...) {
Expand Down Expand Up @@ -4352,8 +4356,8 @@ piEnqueueKernelLaunch(pi_queue Queue, pi_kernel Kernel, pi_uint32 WorkDim,
return PI_SUCCESS;
}

pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context, bool,
pi_kernel *) {
pi_result piextKernelCreateWithNativeHandle(pi_native_handle, pi_context,
pi_kernel *, bool) {
die("Unsupported operation");
return PI_SUCCESS;
}
Expand Down
24 changes: 15 additions & 9 deletions sycl/plugins/level_zero/pi_level_zero.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1056,25 +1056,27 @@ struct _pi_program : _pi_object {
};

// Construct a program in IL or Native state.
_pi_program(pi_context Context, const void *Input, size_t Length, state St)
: State(St), Context(Context), Code(new uint8_t[Length]),
CodeLength(Length), ZeModule(nullptr), HasImports(false),
HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {
_pi_program(pi_context Context, const void *Input, size_t Length, state St,
bool OwnZeModule)
: State(St), Context(Context), OwnZeModule(OwnZeModule),
Code(new uint8_t[Length]), CodeLength(Length), ZeModule(nullptr),
HasImports(false), HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {

std::memcpy(Code.get(), Input, Length);
}

// Construct a program in either Object or Exe state.
_pi_program(pi_context Context, ze_module_handle_t ZeModule, state St,
bool HasImports = false)
: State(St), Context(Context), ZeModule(ZeModule), HasImports(HasImports),
bool OwnZeModule, bool HasImports = false)
: State(St), Context(Context), OwnZeModule(OwnZeModule),
ZeModule(ZeModule), HasImports(HasImports),
HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {}

// Construct a program in LinkedExe state.
_pi_program(pi_context Context, std::vector<LinkedReleaser> &&Inputs,
ze_module_build_log_handle_t ZeLog)
: State(LinkedExe), Context(Context), ZeModule(nullptr),
HasImports(false), HasImportsAndIsLinked(false),
ze_module_build_log_handle_t ZeLog, bool OwnZeModule)
: State(LinkedExe), Context(Context), OwnZeModule(OwnZeModule),
ZeModule(nullptr), HasImports(false), HasImportsAndIsLinked(false),
LinkedPrograms(std::move(Inputs)), ZeBuildLog(ZeLog) {}

~_pi_program();
Expand All @@ -1083,6 +1085,10 @@ struct _pi_program : _pi_object {
state State;
pi_context Context; // Context of the program.

// Indicates if we own the ZeModule or it came from interop that
// asked to not transfer the ownership to SYCL RT.
bool OwnZeModule;

// Used for programs in IL or Native states.
std::unique_ptr<uint8_t[]> Code; // Array containing raw IL / native code.
size_t CodeLength; // Size (bytes) of the array.
Expand Down
10 changes: 6 additions & 4 deletions sycl/plugins/opencl/pi_opencl.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -444,8 +444,9 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length,
}

pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
pi_context,
pi_program *piProgram) {
pi_context, pi_program *piProgram,
bool ownNativeHandle) {
(void)ownNativeHandle;
assert(piProgram != nullptr);
*piProgram = reinterpret_cast<pi_program>(nativeHandle);
return PI_SUCCESS;
Expand Down Expand Up @@ -497,8 +498,9 @@ pi_result piextKernelSetArgSampler(pi_kernel kernel, pi_uint32 arg_index,
}

pi_result piextKernelCreateWithNativeHandle(pi_native_handle nativeHandle,
pi_context, bool,
pi_kernel *piKernel) {
pi_context, pi_kernel *piKernel,
bool ownNativeHandle) {
(void)ownNativeHandle;
assert(piKernel != nullptr);
*piKernel = reinterpret_cast<pi_kernel>(nativeHandle);
return PI_SUCCESS;
Expand Down
16 changes: 14 additions & 2 deletions sycl/source/backend.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -119,12 +119,18 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
std::shared_ptr<detail::kernel_bundle_impl>
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
bundle_state State, backend Backend) {
return make_kernel_bundle(NativeHandle, TargetContext, false, State, Backend);
}

std::shared_ptr<detail::kernel_bundle_impl>
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
bool KeepOwnership, bundle_state State, backend Backend) {
const auto &Plugin = getPlugin(Backend);
const auto &ContextImpl = getSyclObjImpl(TargetContext);

pi::PiProgram PiProgram = nullptr;
Plugin.call<PiApiKind::piextProgramCreateWithNativeHandle>(
NativeHandle, ContextImpl->getHandleRef(), &PiProgram);
NativeHandle, ContextImpl->getHandleRef(), &PiProgram, !KeepOwnership);

std::vector<pi::PiDevice> ProgramDevices;
size_t NumDevices = 0;
Expand Down Expand Up @@ -194,14 +200,20 @@ make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,

return std::make_shared<kernel_bundle_impl>(TargetContext, Devices, DevImg);
}

kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext,
backend Backend) {
return make_kernel(NativeHandle, TargetContext, false, Backend);
}

kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext,
bool KeepOwnership, backend Backend) {
const auto &Plugin = getPlugin(Backend);
const auto &ContextImpl = getSyclObjImpl(TargetContext);
// Create PI kernel first.
pi::PiKernel PiKernel = nullptr;
Plugin.call<PiApiKind::piextKernelCreateWithNativeHandle>(
NativeHandle, ContextImpl->getHandleRef(), false, &PiKernel);
NativeHandle, ContextImpl->getHandleRef(), &PiKernel, !KeepOwnership);

if (Backend == backend::opencl)
Plugin.call<PiApiKind::piKernelRetain>(PiKernel);
Expand Down
Loading