Skip to content

Commit a51c333

Browse files
againullagainull
authored
[SYCL] Implementation of getNative() interoperability for Level Zero (#1723)
* 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 <[email protected]> Signed-off-by: Artur Gainullin <[email protected]>
1 parent dddac4a commit a51c333

22 files changed

+275
-25
lines changed

‎sycl/CMakeLists.txt

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -9,7 +9,7 @@ option(SYCL_ENABLE_WERROR "Treat all warnings as errors in SYCL project" OFF)
99
option(SYCL_ADD_DEV_VERSION_POSTFIX "Adds -V postfix to version string" ON)
1010

1111
set(SYCL_MAJOR_VERSION 2)
12-
set(SYCL_MINOR_VERSION 0)
12+
set(SYCL_MINOR_VERSION 1)
1313
set(SYCL_PATCH_VERSION 0)
1414
set(SYCL_DEV_ABI_VERSION 0)
1515
if (SYCL_ADD_DEV_VERSION_POSTFIX)

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

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -21,6 +21,7 @@ typedef int CUdevice;
2121
typedef struct CUctx_st *CUcontext;
2222
typedef struct CUstream_st *CUstream;
2323
typedef struct CUevent_st *CUevent;
24+
typedef struct CUmod_st *CUmodule;
2425

2526
// As defined in the CUDA 10.1 header file. This requires CUDA version > 3.2
2627
#if defined(_WIN64) || defined(__LP64__)
@@ -40,6 +41,8 @@ template <> struct interop<backend::cuda, queue> { using type = CUstream; };
4041

4142
template <> struct interop<backend::cuda, event> { using type = CUevent; };
4243

44+
template <> struct interop<backend::cuda, program> { using type = CUmodule; };
45+
4346
template <typename DataT, int Dimensions, access::mode AccessMode>
4447
struct interop<backend::cuda, accessor<DataT, Dimensions, AccessMode,
4548
access::target::global_buffer,

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

Lines changed: 81 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,81 @@
1+
//==--------- level_zero.hpp - SYCL Level-Zero backend ---------------------==//
2+
//
3+
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
// See https://llvm.org/LICENSE.txt for license information.
5+
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
//
7+
//===----------------------------------------------------------------------===//
8+
9+
#pragma once
10+
11+
#include <CL/sycl.hpp>
12+
#include <level_zero/ze_api.h>
13+
14+
__SYCL_INLINE_NAMESPACE(cl) {
15+
namespace sycl {
16+
17+
template <> struct interop<backend::level0, platform> {
18+
using type = ze_driver_handle_t;
19+
};
20+
21+
template <> struct interop<backend::level0, device> {
22+
using type = ze_device_handle_t;
23+
};
24+
25+
template <> struct interop<backend::level0, queue> {
26+
using type = ze_command_queue_handle_t;
27+
};
28+
29+
template <> struct interop<backend::level0, program> {
30+
using type = ze_module_handle_t;
31+
};
32+
33+
template <typename DataT, int Dimensions, access::mode AccessMode>
34+
struct interop<backend::level0, accessor<DataT, Dimensions, AccessMode,
35+
access::target::global_buffer,
36+
access::placeholder::false_t>> {
37+
using type = char *;
38+
};
39+
40+
namespace level0 {
41+
42+
// Implementation of various "make" functions resides in libsycl.so
43+
platform make_platform(pi_native_handle NativeHandle);
44+
device make_device(const platform &Platform, pi_native_handle NativeHandle);
45+
program make_program(const context &Context, pi_native_handle NativeHandle);
46+
queue make_queue(const context &Context, pi_native_handle InteropHandle);
47+
48+
// Construction of SYCL platform.
49+
template <typename T, typename std::enable_if<
50+
std::is_same<T, platform>::value>::type * = nullptr>
51+
T make(typename interop<backend::level0, T>::type Interop) {
52+
return make_platform(reinterpret_cast<pi_native_handle>(Interop));
53+
}
54+
55+
// Construction of SYCL device.
56+
template <typename T, typename std::enable_if<
57+
std::is_same<T, device>::value>::type * = nullptr>
58+
T make(const platform &Platform,
59+
typename interop<backend::level0, T>::type Interop) {
60+
return make_device(Platform, reinterpret_cast<pi_native_handle>(Interop));
61+
}
62+
63+
// Construction of SYCL program.
64+
template <typename T, typename std::enable_if<
65+
std::is_same<T, program>::value>::type * = nullptr>
66+
T make(const context &Context,
67+
typename interop<backend::level0, T>::type Interop) {
68+
return make_program(Context, reinterpret_cast<pi_native_handle>(Interop));
69+
}
70+
71+
// Construction of SYCL queue.
72+
template <typename T, typename std::enable_if<
73+
std::is_same<T, queue>::value>::type * = nullptr>
74+
T make(const context &Context,
75+
typename interop<backend::level0, T>::type Interop) {
76+
return make_queue(Context, reinterpret_cast<pi_native_handle>(Interop));
77+
}
78+
79+
} // namespace level0
80+
} // namespace sycl
81+
} // __SYCL_INLINE_NAMESPACE(cl)

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

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -850,9 +850,10 @@ piextDeviceGetNativeHandle(pi_device device, pi_native_handle *nativeHandle);
850850
/// NOTE: The created PI object takes ownership of the native handle.
851851
///
852852
/// \param nativeHandle is the native handle to create PI device from.
853+
/// \param platform is the platform of the device.
853854
/// \param device is the PI device created from the native handle.
854855
__SYCL_EXPORT pi_result piextDeviceCreateWithNativeHandle(
855-
pi_native_handle nativeHandle, pi_device *device);
856+
pi_native_handle nativeHandle, pi_platform platform, pi_device *device);
856857

857858
/// Selects the most appropriate device binary based on runtime information
858859
/// and the IR characteristics.
@@ -944,9 +945,10 @@ piextQueueGetNativeHandle(pi_queue queue, pi_native_handle *nativeHandle);
944945
/// NOTE: The created PI object takes ownership of the native handle.
945946
///
946947
/// \param nativeHandle is the native handle to create PI queue from.
948+
/// \param context is the PI context of the queue.
947949
/// \param queue is the PI queue created from the native handle.
948950
__SYCL_EXPORT pi_result piextQueueCreateWithNativeHandle(
949-
pi_native_handle nativeHandle, pi_queue *queue);
951+
pi_native_handle nativeHandle, pi_context context, pi_queue *queue);
950952

951953
//
952954
// Memory
@@ -1066,9 +1068,10 @@ piextProgramGetNativeHandle(pi_program program, pi_native_handle *nativeHandle);
10661068
/// NOTE: The created PI object takes ownership of the native handle.
10671069
///
10681070
/// \param nativeHandle is the native handle to create PI program from.
1071+
/// \param context is the PI context of the program.
10691072
/// \param program is the PI program created from the native handle.
10701073
__SYCL_EXPORT pi_result piextProgramCreateWithNativeHandle(
1071-
pi_native_handle nativeHandle, pi_program *program);
1074+
pi_native_handle nativeHandle, pi_context context, pi_program *program);
10721075

10731076
//
10741077
// Kernel

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

Lines changed: 3 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -334,14 +334,15 @@ template <class To, class From> inline To cast(From value) {
334334

335335
// These conversions should use PI interop API.
336336
template <> inline pi::PiProgram cast(cl_program) {
337-
RT::assertion(false, "pi::cast -> use piextProgramFromNative");
337+
RT::assertion(false, "pi::cast -> use piextCreateProgramWithNativeHandle");
338338
return {};
339339
}
340340

341341
template <> inline pi::PiDevice cast(cl_device_id) {
342-
RT::assertion(false, "pi::cast -> use piextDeviceFromNative");
342+
RT::assertion(false, "pi::cast -> use piextCreateDeviceWithNativeHandle");
343343
return {};
344344
}
345+
345346
} // namespace pi
346347
} // namespace detail
347348

‎sycl/include/CL/sycl/platform.hpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -112,7 +112,7 @@ class __SYCL_EXPORT platform {
112112
/// \return a native handle, the type of which defined by the backend.
113113
template <backend BackendName>
114114
auto get_native() const -> typename interop<BackendName, platform>::type {
115-
return detail::pi::cast<typename interop<BackendName, platform>::type>(
115+
return reinterpret_cast<typename interop<BackendName, platform>::type>(
116116
getNative());
117117
}
118118

‎sycl/include/CL/sycl/program.hpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -322,7 +322,17 @@ class __SYCL_EXPORT program {
322322
#endif // __SYCL_DEVICE_ONLY__
323323
}
324324

325+
/// Gets the native handle of the SYCL platform.
326+
///
327+
/// \return a native handle, the type of which defined by the backend.
328+
template <backend BackendName>
329+
auto get_native() const -> typename interop<BackendName, program>::type {
330+
return reinterpret_cast<typename interop<BackendName, program>::type>(
331+
getNative());
332+
}
333+
325334
private:
335+
pi_native_handle getNative() const;
326336
program(shared_ptr_class<detail::program_impl> impl);
327337

328338
/// Template-free version of get_kernel.

‎sycl/plugins/cuda/pi_cuda.cpp

Lines changed: 6 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -1338,10 +1338,12 @@ pi_result cuda_piextDeviceGetNativeHandle(pi_device device,
13381338
/// NOTE: The created PI object takes ownership of the native handle.
13391339
///
13401340
/// \param[in] nativeHandle The native handle to create PI device object from.
1341+
/// \param[in] platform is the PI platform of the device.
13411342
/// \param[out] device Set to the PI device object created from native handle.
13421343
///
13431344
/// \return TBD
13441345
pi_result cuda_piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle,
1346+
pi_platform platform,
13451347
pi_device *device) {
13461348
cl::sycl::detail::pi::die(
13471349
"Creation of PI device from native handle not implemented");
@@ -1879,10 +1881,12 @@ pi_result cuda_piextQueueGetNativeHandle(pi_queue queue,
18791881
/// NOTE: The created PI object takes ownership of the native handle.
18801882
///
18811883
/// \param[in] nativeHandle The native handle to create PI queue object from.
1884+
/// \param[in] context is the PI context of the queue.
18821885
/// \param[out] queue Set to the PI queue object created from native handle.
18831886
///
18841887
/// \return TBD
18851888
pi_result cuda_piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle,
1889+
pi_context context,
18861890
pi_queue *queue) {
18871891
cl::sycl::detail::pi::die(
18881892
"Creation of PI queue from native handle not implemented");
@@ -2489,10 +2493,12 @@ pi_result cuda_piextProgramGetNativeHandle(pi_program program,
24892493
/// NOTE: The created PI object takes ownership of the native handle.
24902494
///
24912495
/// \param[in] nativeHandle The native handle to create PI program object from.
2496+
/// \param[in] context The PI context of the program.
24922497
/// \param[out] program Set to the PI program object created from native handle.
24932498
///
24942499
/// \return TBD
24952500
pi_result cuda_piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
2501+
pi_context context,
24962502
pi_program *program) {
24972503
cl::sycl::detail::pi::die(
24982504
"Creation of PI program from native handle not implemented");

‎sycl/plugins/level_zero/pi_level0.cpp

Lines changed: 53 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -1182,10 +1182,16 @@ pi_result piextDeviceGetNativeHandle(pi_device Device,
11821182
}
11831183

11841184
pi_result piextDeviceCreateWithNativeHandle(pi_native_handle NativeHandle,
1185+
pi_platform Platform,
11851186
pi_device *Device) {
1187+
assert(NativeHandle);
1188+
assert(Device);
1189+
assert(Platform);
1190+
11861191
// Create PI device from the given L0 device handle.
1187-
die("piextDeviceCreateWithNativeHandle: not supported");
1188-
return PI_SUCCESS;
1192+
auto ZeDevice = pi_cast<ze_device_handle_t>(NativeHandle);
1193+
*Device = new _pi_device(ZeDevice, Platform);
1194+
return (*Device)->initialize();
11891195
}
11901196

11911197
pi_result piContextCreate(const pi_context_properties *Properties,
@@ -1370,13 +1376,24 @@ pi_result piQueueFinish(pi_queue Queue) {
13701376

13711377
pi_result piextQueueGetNativeHandle(pi_queue Queue,
13721378
pi_native_handle *NativeHandle) {
1373-
die("piextQueueGetNativeHandle: not supported");
1379+
assert(Queue);
1380+
assert(NativeHandle);
1381+
1382+
auto ZeQueue = pi_cast<ze_command_queue_handle_t *>(NativeHandle);
1383+
// Extract the L0 queue handle from the given PI queue
1384+
*ZeQueue = Queue->ZeCommandQueue;
13741385
return PI_SUCCESS;
13751386
}
13761387

13771388
pi_result piextQueueCreateWithNativeHandle(pi_native_handle NativeHandle,
1389+
pi_context Context,
13781390
pi_queue *Queue) {
1379-
die("piextQueueCreateWithNativeHandle: not supported");
1391+
assert(NativeHandle);
1392+
assert(Context);
1393+
assert(Queue);
1394+
1395+
auto ZeQueue = pi_cast<ze_command_queue_handle_t>(NativeHandle);
1396+
*Queue = new _pi_queue(ZeQueue, Context);
13801397
return PI_SUCCESS;
13811398
}
13821399

@@ -1873,13 +1890,43 @@ pi_result piProgramRelease(pi_program Program) {
18731890

18741891
pi_result piextProgramGetNativeHandle(pi_program Program,
18751892
pi_native_handle *NativeHandle) {
1876-
die("piextProgramGetNativeHandle: not supported");
1893+
assert(Program);
1894+
assert(NativeHandle);
1895+
1896+
auto ZeModule = pi_cast<ze_module_handle_t *>(NativeHandle);
1897+
// Extract the L0 module handle from the given PI program
1898+
*ZeModule = Program->ZeModule;
18771899
return PI_SUCCESS;
18781900
}
18791901

18801902
pi_result piextProgramCreateWithNativeHandle(pi_native_handle NativeHandle,
1903+
pi_context Context,
18811904
pi_program *Program) {
1882-
die("piextProgramCreateWithNativeHandle: not supported");
1905+
assert(NativeHandle);
1906+
assert(Context);
1907+
assert(Program);
1908+
1909+
auto ZeModule = pi_cast<ze_module_handle_t>(NativeHandle);
1910+
1911+
// Create PI program from the given L0 module handle.
1912+
//
1913+
// TODO: We don't have the real L0 module descriptor with
1914+
// which it was created, but that's only needed for zeModuleCreate,
1915+
// which we don't expect to be called on the interop program.
1916+
//
1917+
ze_module_desc_t ZeModuleDesc = {};
1918+
ZeModuleDesc.version = ZE_MODULE_DESC_VERSION_CURRENT;
1919+
ZeModuleDesc.format = ZE_MODULE_FORMAT_NATIVE;
1920+
ZeModuleDesc.inputSize = 0;
1921+
ZeModuleDesc.pInputModule = nullptr;
1922+
1923+
try {
1924+
*Program = new _pi_program(ZeModule, ZeModuleDesc, Context);
1925+
} catch (const std::bad_alloc &) {
1926+
return PI_OUT_OF_HOST_MEMORY;
1927+
} catch (...) {
1928+
return PI_ERROR_UNKNOWN;
1929+
}
18831930
return PI_SUCCESS;
18841931
}
18851932

‎sycl/plugins/opencl/pi_opencl.cpp

Lines changed: 3 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -160,7 +160,6 @@ static pi_result USMSetIndirectAccess(pi_kernel kernel) {
160160

161161
extern "C" {
162162

163-
// Example of a PI interface that does not map exactly to an OpenCL one.
164163
pi_result piPlatformsGet(pi_uint32 num_entries, pi_platform *platforms,
165164
pi_uint32 *num_platforms) {
166165
cl_int result = clGetPlatformIDs(cast<cl_uint>(num_entries),
@@ -184,7 +183,6 @@ pi_result piextPlatformCreateWithNativeHandle(pi_native_handle nativeHandle,
184183
return PI_SUCCESS;
185184
}
186185

187-
// Example of a PI interface that does not map exactly to an OpenCL one.
188186
pi_result piDevicesGet(pi_platform platform, pi_device_type device_type,
189187
pi_uint32 num_entries, pi_device *devices,
190188
pi_uint32 *num_devices) {
@@ -274,7 +272,7 @@ pi_result piextDeviceSelectBinary(pi_device device, pi_device_binary *images,
274272
}
275273

276274
pi_result piextDeviceCreateWithNativeHandle(pi_native_handle nativeHandle,
277-
pi_device *piDevice) {
275+
pi_platform, pi_device *piDevice) {
278276
assert(piDevice != nullptr);
279277
*piDevice = reinterpret_cast<pi_device>(nativeHandle);
280278
return PI_SUCCESS;
@@ -321,7 +319,7 @@ pi_result piQueueCreate(pi_context context, pi_device device,
321319
}
322320

323321
pi_result piextQueueCreateWithNativeHandle(pi_native_handle nativeHandle,
324-
pi_queue *piQueue) {
322+
pi_context, pi_queue *piQueue) {
325323
assert(piQueue != nullptr);
326324
*piQueue = reinterpret_cast<pi_queue>(nativeHandle);
327325
return PI_SUCCESS;
@@ -406,6 +404,7 @@ pi_result piProgramCreate(pi_context context, const void *il, size_t length,
406404
}
407405

408406
pi_result piextProgramCreateWithNativeHandle(pi_native_handle nativeHandle,
407+
pi_context,
409408
pi_program *piProgram) {
410409
assert(piProgram != nullptr);
411410
*piProgram = reinterpret_cast<pi_program>(nativeHandle);

‎sycl/source/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -95,6 +95,7 @@ endfunction(add_sycl_rt_library)
9595
set(SYCL_SOURCES
9696
"${sycl_inc_dir}/CL/sycl.hpp"
9797
"backend/opencl.cpp"
98+
"backend/level_zero.cpp"
9899
"detail/accessor_impl.cpp"
99100
"detail/buffer_impl.cpp"
100101
"detail/builtins_common.cpp"

0 commit comments

Comments
 (0)