Skip to content

Commit 7ed63c7

Browse files
smaslov-intelsmaslov-intel
committed
[SYCL] Add ownership control for LeveL-Zero kernel_bundle interop.
Signed-off-by: Sergey V Maslov <sergey.v.maslov@intel.com>
1 parent 9c0508b commit 7ed63c7

File tree

12 files changed

+82
-46
lines changed

12 files changed

+82
-46
lines changed

‎sycl/include/CL/sycl/backend.hpp‎

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -108,9 +108,13 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
108108
backend Backend);
109109
__SYCL_EXPORT kernel make_kernel(pi_native_handle NativeHandle,
110110
const context &TargetContext, backend Backend);
111+
// TODO: Unused. Remove when allowed.
111112
__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
112113
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
113114
bundle_state State, backend Backend);
115+
__SYCL_EXPORT std::shared_ptr<detail::kernel_bundle_impl>
116+
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
117+
bool KeepOwnership, bundle_state State, backend Backend);
114118
} // namespace detail
115119

116120
template <backend Backend>
@@ -221,7 +225,7 @@ make_kernel_bundle(const typename backend_traits<Backend>::template input_type<
221225
std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
222226
detail::make_kernel_bundle(
223227
detail::pi::cast<pi_native_handle>(BackendObject), TargetContext,
224-
State, Backend);
228+
false, State, Backend);
225229
return detail::createSyclObjFromImpl<kernel_bundle<State>>(KBImpl);
226230
}
227231
} // namespace sycl

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

Lines changed: 5 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,7 @@
4040
// changes the API version from 3.5 to 4.6.
4141
// 5.7 Added new context and ownership arguments to
4242
// piextEventCreateWithNativeHandle
43+
// 6.8 Added new ownership argument to piextProgramCreateWithNativeHandle.
4344
//
4445
#include "CL/cl.h"
4546
#define _PI_H_VERSION_MAJOR 5
@@ -1219,9 +1220,12 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle);
12191220
///
12201221
/// \param nativeHandle is the native handle to create PI program from.
12211222
/// \param context is the PI context of the program.
1223+
/// \param ownNativeHandle tells if SYCL RT should assume the ownership of
1224+
/// the native handle, if it can.
12221225
/// \param program is the PI program created from the native handle.
12231226
__SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle(
1224-
pi_native_handle nativeHandle, pi_context context, pi_program *program);
1227+
pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle,
1228+
pi_program *program);
12251229

12261230
//
12271231
// Kernel

‎sycl/include/sycl/ext/oneapi/backend/level_zero.hpp‎

Lines changed: 23 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -101,7 +101,11 @@ template <> struct BackendInput<backend::level_zero, event> {
101101

102102
template <bundle_state State>
103103
struct BackendInput<backend::level_zero, kernel_bundle<State>> {
104-
using type = ze_module_handle_t;
104+
using type = struct {
105+
ze_module_handle_t NativeHandle;
106+
ext::oneapi::level_zero::ownership Ownership{
107+
ext::oneapi::level_zero::ownership::transfer};
108+
};
105109
};
106110

107111
template <bundle_state State>
@@ -249,6 +253,24 @@ event make_event<backend::level_zero>(
249253
BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep);
250254
}
251255

256+
// Specialization of sycl::make_kernel_bundle for Level-Zero backend.
257+
template <>
258+
kernel_bundle<bundle_state::executable>
259+
make_kernel_bundle<backend::ext_oneapi_level_zero, bundle_state::executable>(
260+
const backend_input_t<backend::ext_oneapi_level_zero,
261+
kernel_bundle<bundle_state::executable>>
262+
&BackendObject,
263+
const context &TargetContext) {
264+
std::shared_ptr<detail::kernel_bundle_impl> KBImpl =
265+
detail::make_kernel_bundle(
266+
detail::pi::cast<pi_native_handle>(BackendObject.NativeHandle),
267+
TargetContext,
268+
BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep,
269+
bundle_state::executable, backend::ext_oneapi_level_zero);
270+
return detail::createSyclObjFromImpl<kernel_bundle<bundle_state::executable>>(
271+
KBImpl);
272+
}
273+
252274
// TODO: remove this specialization when generic is changed to call
253275
// .GetNative() instead of .get_native() member of kernel_bundle.
254276
template <>

‎sycl/plugins/cuda/pi_cuda.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3173,7 +3173,7 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program,
31733173
///
31743174
/// \return TBD
31753175
pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
3176-
pi_program *) {
3176+
bool, pi_program *) {
31773177
cl::sycl::detail::pi::die(
31783178
"Creation of PI program from native handle not implemented");
31793179
return {};

‎sycl/plugins/esimd_cpu/pi_esimd_cpu.cpp‎

Lines changed: 9 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -1019,19 +1019,15 @@ pi_result piProgramGetBuildInfo(pi_program, pi_device, cl_program_build_info,
10191019
DIE_NO_IMPLEMENTATION;
10201020
}
10211021

1022-
pi_result piProgramRetain(pi_program) {
1023-
DIE_NO_IMPLEMENTATION;
1024-
}
1022+
pi_result piProgramRetain(pi_program) { DIE_NO_IMPLEMENTATION; }
10251023

1026-
pi_result piProgramRelease(pi_program) {
1027-
DIE_NO_IMPLEMENTATION;
1028-
}
1024+
pi_result piProgramRelease(pi_program) { DIE_NO_IMPLEMENTATION; }
10291025

10301026
pi_result piextProgramGetNativeHandle(pi_program, pi_native_handle *) {
10311027
DIE_NO_IMPLEMENTATION;
10321028
}
10331029

1034-
pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context,
1030+
pi_result piextProgramCreateWithNativeHandle(pi_native_handle, pi_context, bool,
10351031
pi_program *) {
10361032
DIE_NO_IMPLEMENTATION;
10371033
}
@@ -1068,17 +1064,11 @@ pi_result piKernelGetSubGroupInfo(pi_kernel, pi_device,
10681064
DIE_NO_IMPLEMENTATION;
10691065
}
10701066

1071-
pi_result piKernelRetain(pi_kernel) {
1072-
DIE_NO_IMPLEMENTATION;
1073-
}
1067+
pi_result piKernelRetain(pi_kernel) { DIE_NO_IMPLEMENTATION; }
10741068

1075-
pi_result piKernelRelease(pi_kernel) {
1076-
DIE_NO_IMPLEMENTATION;
1077-
}
1069+
pi_result piKernelRelease(pi_kernel) { DIE_NO_IMPLEMENTATION; }
10781070

1079-
pi_result piEventCreate(pi_context, pi_event *) {
1080-
DIE_NO_IMPLEMENTATION;
1081-
}
1071+
pi_result piEventCreate(pi_context, pi_event *) { DIE_NO_IMPLEMENTATION; }
10821072

10831073
pi_result piEventGetInfo(pi_event, pi_event_info, size_t, void *, size_t *) {
10841074
DIE_NO_IMPLEMENTATION;
@@ -1117,9 +1107,7 @@ pi_result piEventSetCallback(pi_event, pi_int32,
11171107
DIE_NO_IMPLEMENTATION;
11181108
}
11191109

1120-
pi_result piEventSetStatus(pi_event, pi_int32) {
1121-
DIE_NO_IMPLEMENTATION;
1122-
}
1110+
pi_result piEventSetStatus(pi_event, pi_int32) { DIE_NO_IMPLEMENTATION; }
11231111

11241112
pi_result piEventRetain(pi_event Event) {
11251113
if (Event == nullptr) {
@@ -1170,13 +1158,9 @@ pi_result piSamplerGetInfo(pi_sampler, pi_sampler_info, size_t, void *,
11701158
DIE_NO_IMPLEMENTATION;
11711159
}
11721160

1173-
pi_result piSamplerRetain(pi_sampler) {
1174-
DIE_NO_IMPLEMENTATION;
1175-
}
1161+
pi_result piSamplerRetain(pi_sampler) { DIE_NO_IMPLEMENTATION; }
11761162

1177-
pi_result piSamplerRelease(pi_sampler) {
1178-
DIE_NO_IMPLEMENTATION;
1179-
}
1163+
pi_result piSamplerRelease(pi_sampler) { DIE_NO_IMPLEMENTATION; }
11801164

11811165
pi_result piEnqueueEventsWait(pi_queue, pi_uint32, const pi_event *,
11821166
pi_event *) {

‎sycl/plugins/hip/pi_hip.cpp‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -2949,11 +2949,14 @@ pi_result hip_piextProgramGetNativeHandle(pi_program program,
29492949
///
29502950
/// \param[in] nativeHandle The native handle to create PI program object from.
29512951
/// \param[in] context The PI context of the program.
2952+
/// \param[in] ownNativeHandle tells if should assume the ownership of
2953+
/// the native handle.
29522954
/// \param[out] program Set to the PI program object created from native handle.
29532955
///
29542956
/// \return TBD
29552957
pi_result hip_piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
29562958
pi_context context,
2959+
bool ownNativeHandle,
29572960
pi_program *program) {
29582961
cl::sycl::detail::pi::die(
29592962
"Creation of PI program from native handle not implemented");

‎sycl/plugins/level_zero/pi_level_zero.cpp‎

Lines changed: 7 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -3657,8 +3657,9 @@ pi_result piProgramLink(pi_context Context, pi_uint32 NumDevices,
36573657
if (res != PI_SUCCESS) {
36583658
return res;
36593659
}
3660-
Input = new _pi_program(Input->Context, ZeModule, _pi_program::Object,
3661-
Input->HasImports);
3660+
Input =
3661+
new _pi_program(Input->Context, ZeModule, true /*own ZeModule*/,
3662+
_pi_program::Object, Input->HasImports);
36623663
Input->HasImportsAndIsLinked = true;
36633664
}
36643665
} else {
@@ -3913,6 +3914,7 @@ pi_result piextProgramGetNativeHandle(pi_program Program,
39133914

39143915
pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
39153916
pi_context Context,
3917+
bool ownNativeHandle,
39163918
pi_program *Program) {
39173919
PI_ASSERT(Program, PI_INVALID_PROGRAM);
39183920
PI_ASSERT(NativeHandle, PI_INVALID_VALUE);
@@ -3925,7 +3927,8 @@ pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
39253927
// executable (state Object).
39263928

39273929
try {
3928-
*Program = new _pi_program(Context, ZeModule, _pi_program::Exe);
3930+
*Program =
3931+
new _pi_program(Context, ZeModule, ownNativeHandle, _pi_program::Exe);
39293932
} catch (const std::bad_alloc &) {
39303933
return PI_OUT_OF_HOST_MEMORY;
39313934
} catch (...) {
@@ -3942,7 +3945,7 @@ _pi_program::~_pi_program() {
39423945
ZE_CALL_NOCHECK(zeModuleBuildLogDestroy, (ZeBuildLog));
39433946
}
39443947

3945-
if (ZeModule) {
3948+
if (ZeModule && OwnZeModule) {
39463949
ZE_CALL_NOCHECK(zeModuleDestroy, (ZeModule));
39473950
}
39483951
}

‎sycl/plugins/level_zero/pi_level_zero.hpp‎

Lines changed: 14 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -1068,23 +1068,24 @@ struct _pi_program : _pi_object {
10681068
// Construct a program in IL or Native state.
10691069
_pi_program(pi_context Context, const void *Input, size_t Length, state St)
10701070
: State(St), Context(Context), Code(new uint8_t[Length]),
1071-
CodeLength(Length), ZeModule(nullptr), HasImports(false),
1072-
HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {
1071+
CodeLength(Length), ZeModule(nullptr), OwnZeModule{true},
1072+
HasImports(false), HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {
10731073

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

10771077
// Construct a program in either Object or Exe state.
1078-
_pi_program(pi_context Context, ze_module_handle_t ZeModule, state St,
1079-
bool HasImports = false)
1080-
: State(St), Context(Context), ZeModule(ZeModule), HasImports(HasImports),
1078+
_pi_program(pi_context Context, ze_module_handle_t ZeModule, bool OwnZeModule,
1079+
state St, bool HasImports = false)
1080+
: State(St), Context(Context),
1081+
ZeModule(ZeModule), OwnZeModule{OwnZeModule}, HasImports(HasImports),
10811082
HasImportsAndIsLinked(false), ZeBuildLog(nullptr) {}
10821083

10831084
// Construct a program in LinkedExe state.
10841085
_pi_program(pi_context Context, std::vector<LinkedReleaser> &&Inputs,
10851086
ze_module_build_log_handle_t ZeLog)
10861087
: State(LinkedExe), Context(Context), ZeModule(nullptr),
1087-
HasImports(false), HasImportsAndIsLinked(false),
1088+
OwnZeModule(true), HasImports(false), HasImportsAndIsLinked(false),
10881089
LinkedPrograms(std::move(Inputs)), ZeBuildLog(ZeLog) {}
10891090

10901091
~_pi_program();
@@ -1103,7 +1104,13 @@ struct _pi_program : _pi_object {
11031104

11041105
// Used for programs in Object or Exe state.
11051106
ze_module_handle_t ZeModule; // Level Zero module handle.
1106-
bool HasImports; // Tells if module imports any symbols.
1107+
1108+
// Indicates if we own the ZeModule or it came from interop that
1109+
// asked to not transfer the ownership to SYCL RT.
1110+
bool OwnZeModule;
1111+
1112+
// Tells if module imports any symbols.
1113+
bool HasImports;
11071114

11081115
// Used for programs in Object state. Tells if this module imports any
11091116
// symbols AND it is linked into some other program that has state LinkedExe.

‎sycl/plugins/opencl/pi_opencl.cpp‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -444,7 +444,7 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length,
444444
}
445445

446446
pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
447-
pi_context,
447+
pi_context, bool,
448448
pi_program *piProgram) {
449449
assert(piProgram != nullptr);
450450
*piProgram = reinterpret_cast<pi_program>(nativeHandle);

‎sycl/source/backend.cpp‎

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -118,13 +118,13 @@ __SYCL_EXPORT event make_event(pi_native_handle NativeHandle,
118118

119119
std::shared_ptr<detail::kernel_bundle_impl>
120120
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
121-
bundle_state State, backend Backend) {
121+
bool KeepOwnership, bundle_state State, backend Backend) {
122122
const auto &Plugin = getPlugin(Backend);
123123
const auto &ContextImpl = getSyclObjImpl(TargetContext);
124124

125125
pi::PiProgram PiProgram = nullptr;
126126
Plugin.call<PiApiKind::piextProgramCreateWithNativeHandle>(
127-
NativeHandle, ContextImpl->getHandleRef(), &PiProgram);
127+
NativeHandle, ContextImpl->getHandleRef(), KeepOwnership, &PiProgram);
128128

129129
std::vector<pi::PiDevice> ProgramDevices;
130130
size_t NumDevices = 0;
@@ -194,6 +194,14 @@ make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
194194

195195
return std::make_shared<kernel_bundle_impl>(TargetContext, Devices, DevImg);
196196
}
197+
198+
// TODO: Unused. Remove when allowed.
199+
std::shared_ptr<detail::kernel_bundle_impl>
200+
make_kernel_bundle(pi_native_handle NativeHandle, const context &TargetContext,
201+
bundle_state State, backend Backend) {
202+
return make_kernel_bundle(NativeHandle, TargetContext, false, State, Backend);
203+
}
204+
197205
kernel make_kernel(pi_native_handle NativeHandle, const context &TargetContext,
198206
backend Backend) {
199207
const auto &Plugin = getPlugin(Backend);

0 commit comments

Comments
 (0)