From cb078faa6ae9fcf0f0c3587d96d723e700b0bf63 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 9 Mar 2023 20:07:42 -0800 Subject: [PATCH 01/19] New interop support for images for LevelZero. Includes make_image and interop_handle::get_native_mem. Signed-off-by: Chris Perkins --- .../sycl_ext_oneapi_backend_level_zero.md | 50 +++++++++- sycl/include/sycl/backend.hpp | 15 +++ sycl/include/sycl/detail/backend_traits.hpp | 2 + .../sycl/detail/backend_traits_level_zero.hpp | 22 +++++ .../sycl/detail/backend_traits_opencl.hpp | 1 + sycl/include/sycl/detail/pi.def | 1 + sycl/include/sycl/detail/pi.h | 17 +++- .../sycl/ext/oneapi/backend/level_zero.hpp | 18 ++++ .../backend/backend_traits_cuda.hpp | 1 + sycl/include/sycl/image.hpp | 43 ++++++++ sycl/include/sycl/interop_handle.hpp | 41 +++++++- sycl/plugins/cuda/pi_cuda.cpp | 20 ++++ .../esimd_emulator/pi_esimd_emulator.cpp | 6 ++ sycl/plugins/hip/pi_hip.cpp | 29 ++++++ sycl/plugins/level_zero/pi_level_zero.cpp | 97 +++++++++++++++---- sycl/plugins/level_zero/pi_level_zero.hpp | 9 +- sycl/plugins/opencl/pi_opencl.cpp | 15 +++ sycl/source/detail/image_impl.cpp | 18 ++++ sycl/source/detail/image_impl.hpp | 6 ++ sycl/source/detail/sycl_mem_obj_t.cpp | 59 +++++++++++ sycl/source/detail/sycl_mem_obj_t.hpp | 6 ++ sycl/source/image.cpp | 11 +++ sycl/unittests/helpers/PiMockPlugin.hpp | 9 ++ 23 files changed, 465 insertions(+), 31 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index d115a58777eb6..c4100b6b39f51 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -15,6 +15,7 @@ This extension provides a feature-test macro as described in the core SYCL speci |1|Initial extension version. |2|Added support for the make_buffer() API. |3|Added device member to backend_input_t. +|4|Added support for make_image() API. NOTE: This extension is following SYCL 2020 backend specification. Prior API for interoperability with Level-Zero is marked as deprecated and will be removed in the next release. @@ -40,15 +41,15 @@ There are multiple ways in which the Level-Zero backend can be selected by the u ### 3.1 Through an environment variable -The SYCL_DEVICE_FILTER environment variable limits the SYCL runtime to use only a subset of the system's devices. -By using ```level_zero``` for backend in SYCL_DEVICE_FILTER you can select the use of Level-Zero as a SYCL backend. +The ONEAPI_DEVICE_SELECTOR environment variable limits the SYCL runtime to use only a subset of the system's devices. +By using ```level_zero``` for backend in ONEAPI_DEVICE_SELECTOR you can select the use of Level-Zero as a SYCL backend. For further details see here: . ### 3.2 Through a programming API There is an extension that introduces a filtering device selection to SYCL described in [sycl\_ext\_oneapi\_filter\_selector](../supported/sycl_ext_oneapi_filter_selector.asciidoc). -Similar to how SYCL_DEVICE_FILTER applies filtering to the entire process this device selector can be used to +Similar to how SYCL_DEVICE_FILTER or ONEAPI_DEVICE_SELECTOR applies filtering to the entire process this device selector can be used to programmatically select the Level-Zero backend. When neither the environment variable nor the filtering device selector are used, the implementation chooses @@ -206,6 +207,28 @@ struct { ``` + +image + + +``` C++ +ze_image_handle_t +``` + + + +``` C++ +struct { + ze_image_handle_t ZeImageHandle; + sycl::image_channel_order ChanOrder; + sycl::image_channel_type ChanType; + range Range; + ext::oneapi::level_zero::ownership Ownership{ + ext::oneapi::level_zero::ownership::transfer}; + } +``` + + [^1]: The SYCL implementation is responsible for distinguishing between the variants of backend_input_t. @@ -223,7 +246,7 @@ It is currently supported for SYCL ```platform```, ```device```, ```context```, ```kernel_bundle```, and ```kernel``` classes. The ```sycl::get_native``` -free-function is not supported for SYCL ```buffer``` class. The native backend object associated with the +free-function is not supported for SYCL ```buffer``` or ```image``` class. The native backend object associated with the buffer can be obtained using interop_hande class as described in the core SYCL specification section 4.10.2, "Class interop_handle". The pointer returned by ```get_native_mem``` method of the ```interop_handle``` @@ -389,6 +412,24 @@ Construct a SYCL buffer instance from a pointer to a Level Zero memory allocatio description above for semantics and restrictions. The additional AvailableEvent argument must be a valid SYCL event. The instance of the SYCL buffer class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used. + + + + +``` C++ +make_image( + const backend_input_t> &, + const context &Context, event AvailableEvent = {}) +``` + +This API is available starting with revision 4 of this specification. + +Construct a SYCL image instance from a ze_image_handle_t. The input SYCL context Context must be associated with a single device, matching the device used at the prior allocation. +The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. +The Ownership input structure member 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. If the behavior is "transfer" then the runtime is going to free the input Level-Zero memory allocation. +Synchronization rules for a buffer that is created with this API are described in Section 4.5 + NOTE: We shall consider adding other interoperability as needed, if possible. @@ -465,3 +506,4 @@ The behavior of the SYCL buffer destructor depends on the Ownership flag. As wit |8|2022-01-06|Artur Gainullin|Introduced make_buffer() API |9|2022-05-12|Steffen Larsen|Added device member to queue input type |10|2022-08-18|Sergey Maslov|Moved free_memory device info query to be sycl_ext_intel_device_info extension +|11|2023-03-08|Chris Perkins|Introduced make_image() API diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp index b12b0310206f2..fa2b8d5c054ff 100644 --- a/sycl/include/sycl/backend.hpp +++ b/sycl/include/sycl/backend.hpp @@ -14,6 +14,7 @@ #include #include #include +#include #if SYCL_BACKEND_OPENCL #include #endif @@ -309,6 +310,20 @@ make_buffer(const typename backend_traits::template input_type< AvailableEvent); } +template +typename std::enable_if::MakeImage == + true && + Backend != backend::ext_oneapi_level_zero, + image>::type +make_image(const typename backend_traits::template input_type< + image> &BackendObject, + const context &TargetContext, event AvailableEvent = {}) { + return image( + detail::pi::cast(BackendObject), TargetContext, + AvailableEvent); +} + template kernel make_kernel(const typename backend_traits::template input_type diff --git a/sycl/include/sycl/detail/backend_traits.hpp b/sycl/include/sycl/detail/backend_traits.hpp index 3317b5632dcb4..7f9cf23ec325a 100644 --- a/sycl/include/sycl/detail/backend_traits.hpp +++ b/sycl/include/sycl/detail/backend_traits.hpp @@ -29,6 +29,8 @@ template struct InteropFeatureSupportMap { static constexpr bool MakeEvent = false; static constexpr bool MakeBuffer = false; static constexpr bool MakeKernel = false; + static constexpr bool MakeKernelBundle = false; + static constexpr bool MakeImage = false; }; } // namespace detail } // __SYCL_INLINE_VER_NAMESPACE(_V1) diff --git a/sycl/include/sycl/detail/backend_traits_level_zero.hpp b/sycl/include/sycl/detail/backend_traits_level_zero.hpp index add23a8cda27f..50f1ac875a2a9 100644 --- a/sycl/include/sycl/detail/backend_traits_level_zero.hpp +++ b/sycl/include/sycl/detail/backend_traits_level_zero.hpp @@ -159,6 +159,27 @@ struct BackendReturn +struct BackendInput> { + // LevelZero has no way of getting image description FROM a ZeImageHandle so + // it must be provided. + struct type { + ze_image_handle_t ZeImageHandle; + sycl::image_channel_order ChanOrder; + sycl::image_channel_type ChanType; + range Range; + ext::oneapi::level_zero::ownership Ownership{ + ext::oneapi::level_zero::ownership::transfer}; + }; +}; + +template +struct BackendReturn> { + using type = ze_image_handle_t; +}; + template <> struct BackendReturn { using type = ze_command_queue_handle_t; }; @@ -207,6 +228,7 @@ template <> struct InteropFeatureSupportMap { static constexpr bool MakeKernelBundle = true; static constexpr bool MakeKernel = true; static constexpr bool MakeBuffer = true; + static constexpr bool MakeImage = true; }; } // namespace detail diff --git a/sycl/include/sycl/detail/backend_traits_opencl.hpp b/sycl/include/sycl/detail/backend_traits_opencl.hpp index b943389a0c29a..494cc725b61ed 100644 --- a/sycl/include/sycl/detail/backend_traits_opencl.hpp +++ b/sycl/include/sycl/detail/backend_traits_opencl.hpp @@ -155,6 +155,7 @@ template <> struct InteropFeatureSupportMap { static constexpr bool MakeBuffer = true; static constexpr bool MakeKernel = true; static constexpr bool MakeKernelBundle = true; + static constexpr bool MakeImage = false; }; namespace pi { diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index eda09035c883e..11cf51f86be30 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -61,6 +61,7 @@ _PI_API(piMemRelease) _PI_API(piMemBufferPartition) _PI_API(piextMemGetNativeHandle) _PI_API(piextMemCreateWithNativeHandle) +_PI_API(piextImgCreateWithNativeHandle) // Program _PI_API(piProgramCreate) _PI_API(piclProgramCreateWithSource) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 56b8b33fae583..ee9c126944add 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -77,9 +77,11 @@ // 12.22 Add piGetDeviceAndHostTimer to query device wall-clock timestamp // 12.23 Added new piextEnqueueDeviceGlobalVariableWrite and // piextEnqueueDeviceGlobalVariableRead functions. +// 12.24 Added piextImgCreateWithNativeHandle for creating images from native +// handles. #define _PI_H_VERSION_MAJOR 12 -#define _PI_H_VERSION_MINOR 23 +#define _PI_H_VERSION_MINOR 24 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) @@ -1249,6 +1251,19 @@ __SYCL_EXPORT pi_result piextMemCreateWithNativeHandle( pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, pi_mem *mem); +/// Creates PI image object from a native handle. +/// NOTE: The created PI object takes ownership of the native handle. +/// +/// \param nativeHandle is the native handle to create PI image from. +/// \param context The PI context of the memory allocation. +/// \param ownNativeHandle Indicates if we own the native memory handle or it +/// came from interop that asked to not transfer the ownership to SYCL RT. +/// \param img is the PI img created from the native handle. +__SYCL_EXPORT pi_result piextImgCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *img); + // // Program // diff --git a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp index 6c3e1b63e70ff..1dec8a876e5a6 100644 --- a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp +++ b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp @@ -198,6 +198,24 @@ make_buffer( !(BackendObject.Ownership == ext::oneapi::level_zero::ownership::keep)); } +// Specialization of sycl::make_image for Level-Zero backend. +template +typename std::enable_if>::type +make_image(const backend_input_t> + &BackendObject, + const context &TargetContext, event AvailableEvent) { + + bool OwnNativeHandle = + (BackendObject.Ownership == ext::oneapi::level_zero::ownership::transfer); + + return image( + detail::pi::cast(BackendObject.ZeImageHandle), + TargetContext, AvailableEvent, BackendObject.ChanOrder, + BackendObject.ChanType, OwnNativeHandle, BackendObject.Range); +} + namespace __SYCL2020_DEPRECATED("use 'ext::oneapi::level_zero' instead") level_zero { using namespace ext::oneapi::level_zero; diff --git a/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp b/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp index 370866eb126d5..93b8c760e1081 100644 --- a/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp +++ b/sycl/include/sycl/ext/oneapi/experimental/backend/backend_traits_cuda.hpp @@ -120,6 +120,7 @@ template <> struct InteropFeatureSupportMap { static constexpr bool MakeBuffer = false; static constexpr bool MakeKernel = false; static constexpr bool MakeKernelBundle = false; + static constexpr bool MakeImage = false; }; } // namespace detail diff --git a/sycl/include/sycl/image.hpp b/sycl/include/sycl/image.hpp index 873f51068a60e..1bc001ae0f577 100644 --- a/sycl/include/sycl/image.hpp +++ b/sycl/include/sycl/image.hpp @@ -22,8 +22,18 @@ namespace sycl { __SYCL_INLINE_VER_NAMESPACE(_V1) { +// forward declarations class handler; +template class image; + +// 'friend' +template +typename std::enable_if>::type +make_image(const backend_input_t> &BackendObject, + const context &TargetContext, event AvailableEvent = {}); + enum class image_channel_order : unsigned int { a = 0, r = 1, @@ -128,6 +138,13 @@ class __SYCL_EXPORT image_plain { uint8_t Dimensions); #endif + image_plain(pi_native_handle MemObject, const context &SyclContext, + event AvailableEvent, + std::unique_ptr Allocator, + uint8_t Dimensions, image_channel_order Order, + image_channel_type Type, bool OwnNativeHandle, + range<3> Range3WithOnes); + template bool has_property() const noexcept; template propertyT get_property() const; @@ -467,6 +484,15 @@ class image : public detail::image_plain { void set_write_back(bool flag = true) { image_plain::set_write_back(flag); } private: + image(pi_native_handle MemObject, const context &SyclContext, + event AvailableEvent, image_channel_order Order, + image_channel_type Type, bool OwnNativeHandle, range Range) + : image_plain(MemObject, SyclContext, AvailableEvent, + make_unique_ptr< + detail::SYCLMemObjAllocatorHolder>(), + Dimensions, Order, Type, OwnNativeHandle, + detail::convertToArrayOfN<3, 1>(Range)) {} + // This utility api is currently used by accessor to get the element size of // the image. Element size is dependent on num of channels and channel type. // This information is not accessible from the image using any public API. @@ -484,6 +510,23 @@ class image : public detail::image_plain { return image_plain::getChannelType(); } + // Declare make_image as a friend function + template + friend typename std::enable_if< + detail::InteropFeatureSupportMap::MakeImage == true && + Backend != backend::ext_oneapi_level_zero, + image>::type + make_image( + const typename backend_traits::template input_type> + &BackendObject, + const context &TargetContext, event AvailableEvent); + + template + friend typename std::enable_if>::type + make_image(const backend_input_t> &BackendObject, + const context &TargetContext, event AvailableEvent); + template friend decltype(Obj::impl) detail::getSyclObjImpl(const Obj &SyclObject); diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index 8804073f827a9..b045b8b766597 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -50,9 +50,10 @@ class interop_handle { template > - backend_return_t> get_native_mem( - const accessor &Acc) - const { + detail::enable_if_t>> + get_native_mem(const accessor + &Acc) const { static_assert(Target == access::target::device || Target == access::target::constant_buffer, "The method is available only for target::device accessors"); @@ -70,6 +71,33 @@ class interop_handle { #endif } + /// Receives a SYCL accessor that has been defined as a requirement for the + /// command group, and returns the underlying OpenCL memory object that is + /// used by the SYCL runtime. If the accessor passed as parameter is not part + /// of the command group requirements (e.g. it is an unregistered placeholder + /// accessor), the exception `sycl::invalid_object` is thrown + /// asynchronously. + template //, + // typename PropertyListT = ext::oneapi::accessor_property_list<>> + backend_return_t> get_native_mem( + const detail::image_accessor &Acc) const { +#ifndef __SYCL_DEVICE_ONLY__ + if (Backend != get_backend()) + throw invalid_object_error("Incorrect backend argument was passed", + PI_ERROR_INVALID_MEM_OBJECT); + const auto *AccBase = static_cast(&Acc); + return getMemImpl(detail::getSyclObjImpl(*AccBase).get()); +#else + (void)Acc; + // we believe this won't be ever called on device side + return backend_return_t>{0}; +#endif + } + /// Returns an underlying native backend object associated with teh queue /// that the host task was submitted to. If the command group was submitted /// with a secondary queue and the fall-back was triggered, the queue that @@ -162,6 +190,13 @@ class interop_handle { NativeHandles); } + template + backend_return_t> + getMemImpl(detail::AccessorImplHost *Req) const { + using image_return_t = backend_return_t>; + return reinterpret_cast(getNativeMem(Req)); + } + __SYCL_EXPORT pi_native_handle getNativeMem(detail::AccessorImplHost *Req) const; __SYCL_EXPORT pi_native_handle getNativeQueue() const; diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 2d88978d87780..fe56b4bd3bda5 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2433,6 +2433,26 @@ pi_result cuda_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return {}; } +/// Created a PI image mem object from a CUDA image mem handle. +/// TODO: Implement this. +/// NOTE: The created PI object takes ownership of the native handle. +/// +/// \param[in] nativeHandle The native handle to create PI mem object from. +/// \param[in] context The PI context of the memory allocation. +/// \param[in] ownNativeHandle Indicates if we own the native memory handle or +/// it came from interop that asked to not transfer the ownership to SYCL RT. +/// \param[out] mem Set to the PI mem object created from native handle. +/// +/// \return TBD +pi_result cuda_piextImgCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *Img) { + sycl::detail::pi::die( + "Creation of PI mem from native image handle not implemented"); + return {}; +} + /// Creates a `pi_queue` object on the CUDA backend. /// Valid properties /// * __SYCL_PI_CUDA_USE_DEFAULT_STREAM -> CU_STREAM_DEFAULT diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 0fc2a5a10f4f9..93ffc7c5c3ae2 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1297,6 +1297,12 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, DIE_NO_IMPLEMENTATION; } +pi_result piextImgCreateWithNativeHandle(pi_native_handle, pi_context, bool, + const pi_image_format *, + const pi_image_desc *, pi_mem *) { + DIE_NO_IMPLEMENTATION; +} + pi_result piProgramCreate(pi_context, const void *, size_t, pi_program *) { DIE_NO_IMPLEMENTATION; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index ddca2a872adfc..dd0e31a938f67 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -2382,6 +2382,35 @@ pi_result hip_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return {}; } +/// Created a PI image mem object from a HIP image mem handle. +/// TODO: Implement this. +/// NOTE: The created PI object takes ownership of the native handle. +/// +/// \param[in] nativeHandle The native handle to create PI mem object from. +/// \param[in] context The PI context of the memory allocation. +/// \param[in] ownNativeHandle Indicates if we own the native memory handle or +/// it came from interop that asked to not transfer the ownership to SYCL RT. +/// \param[out] mem Set to the PI mem object created from native handle. +/// +/// \return TBD +pi_result hip_piextImgCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + bool ownNativeHandle, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + pi_mem *mem) { + (void)nativeHandle; + (void)context; + (void)ownNativeHandle; + (void)ImageFormat; + (void)ImageDesc; + (void)mem; + + sycl::detail::pi::die( + "Creation of PI mem from native image handle not implemented"); + return {}; +} + /// Creates a `pi_queue` object on the HIP backend. /// Valid properties /// * __SYCL_PI_HIP_USE_DEFAULT_STREAM -> hipStreamDefault diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 984bd342989bd..e4fda0fa9a915 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3015,8 +3015,10 @@ pi_result piMemBufferCreate(pi_context Context, pi_mem_flags Flags, size_t Size, pi_result piMemGetInfo(pi_mem Mem, pi_mem_info ParamName, size_t ParamValueSize, void *ParamValue, size_t *ParamValueSizeRet) { PI_ASSERT(Mem, PI_ERROR_INVALID_VALUE); - // piMemImageGetInfo must be used for images - PI_ASSERT(!Mem->isImage(), PI_ERROR_INVALID_VALUE); + // piMemImageGetInfo must be used for images, except for shared params (like + // Context, AccessMode, etc) + PI_ASSERT(ParamName == PI_MEM_CONTEXT || !Mem->isImage(), + PI_ERROR_INVALID_VALUE); std::shared_lock Lock(Mem->Mutex); ReturnHelper ReturnValue(ParamValueSize, ParamValue, ParamValueSizeRet); @@ -3087,8 +3089,11 @@ pi_result piMemRelease(pi_mem Mem) { if (Mem->isImage()) { char *ZeHandleImage; - PI_CALL(Mem->getZeHandle(ZeHandleImage, _pi_mem::write_only)); - ZE_CALL(zeImageDestroy, (pi_cast(ZeHandleImage))); + auto Image = static_cast(Mem); + if (Image->OwnZeMemHandle) { + PI_CALL(Mem->getZeHandle(ZeHandleImage, _pi_mem::write_only)); + ZE_CALL(zeImageDestroy, (pi_cast(ZeHandleImage))); + } } else { auto Buffer = static_cast(Mem); Buffer->free(); @@ -3098,20 +3103,9 @@ pi_result piMemRelease(pi_mem Mem) { return PI_SUCCESS; } -pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, void *HostPtr, - pi_mem *RetImage) { - - // TODO: implement read-only, write-only - if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) { - die("piMemImageCreate: Level-Zero implements only read-write buffer," - "no read-only or write-only yet."); - } - PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT); - PI_ASSERT(RetImage, PI_ERROR_INVALID_VALUE); - PI_ASSERT(ImageFormat, PI_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR); - +pi_result PIToZeImageDesc(const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + ZeStruct &ZeImageDesc) { ze_image_format_type_t ZeImageFormatType; size_t ZeImageFormatTypeSize; switch (ImageFormat->image_channel_data_type) { @@ -3222,8 +3216,8 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, return PI_ERROR_INVALID_VALUE; } - ZeStruct ZeImageDesc; - ZeImageDesc.arraylevels = ZeImageDesc.flags = 0; + ZeImageDesc.arraylevels = 0; + ZeImageDesc.flags = 0; ZeImageDesc.type = ZeImageType; ZeImageDesc.format = ZeFormatDesc; ZeImageDesc.width = pi_cast(ImageDesc->image_width); @@ -3232,6 +3226,29 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, ZeImageDesc.arraylevels = pi_cast(ImageDesc->image_array_size); ZeImageDesc.miplevels = ImageDesc->num_mip_levels; + return PI_SUCCESS; +} + +pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, void *HostPtr, + pi_mem *RetImage) { + + // TODO: implement read-only, write-only + if ((Flags & PI_MEM_FLAGS_ACCESS_RW) == 0) { + die("piMemImageCreate: Level-Zero implements only read-write buffer," + "no read-only or write-only yet."); + } + PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT); + PI_ASSERT(RetImage, PI_ERROR_INVALID_VALUE); + PI_ASSERT(ImageFormat, PI_ERROR_INVALID_IMAGE_FORMAT_DESCRIPTOR); + + ZeStruct ZeImageDesc; + pi_result DescriptionResult = + PIToZeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); + if (DescriptionResult != PI_SUCCESS) + return DescriptionResult; + std::shared_lock Lock(Context->Mutex); // Currently we have the "0" device in context with mutliple root devices to @@ -3245,7 +3262,7 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, (Context->ZeContext, Device->ZeDevice, &ZeImageDesc, &ZeHImage)); try { - auto ZePIImage = new _pi_image(Context, ZeHImage); + auto ZePIImage = new _pi_image(Context, ZeHImage, /*OwnNativeHandle=*/true); *RetImage = ZePIImage; #ifndef NDEBUG @@ -3370,6 +3387,44 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle NativeHandle, return PI_SUCCESS; } +pi_result piextImgCreateWithNativeHandle(pi_native_handle NativeHandle, + pi_context Context, + bool OwnNativeHandle, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + pi_mem *RetImage) { + + PI_ASSERT(RetImage, PI_ERROR_INVALID_VALUE); + PI_ASSERT(NativeHandle, PI_ERROR_INVALID_VALUE); + PI_ASSERT(Context, PI_ERROR_INVALID_CONTEXT); + + std::shared_lock Lock(Context->Mutex); + + ze_image_handle_t ZeHImage = pi_cast(NativeHandle); + + try { + auto ZePIImage = new _pi_image(Context, ZeHImage, OwnNativeHandle); + *RetImage = ZePIImage; + +#ifndef NDEBUG + ZeStruct ZeImageDesc; + pi_result DescriptionResult = + PIToZeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); + if (DescriptionResult != PI_SUCCESS) + return DescriptionResult; + + ZePIImage->ZeImageDesc = ZeImageDesc; +#endif // !NDEBUG + + } catch (const std::bad_alloc &) { + return PI_ERROR_OUT_OF_HOST_MEMORY; + } catch (...) { + return PI_ERROR_UNKNOWN; + } + + return PI_SUCCESS; +} + pi_result piProgramCreate(pi_context Context, const void *ILBytes, size_t Length, pi_program *Program) { diff --git a/sycl/plugins/level_zero/pi_level_zero.hpp b/sycl/plugins/level_zero/pi_level_zero.hpp index f4b7d1098efe5..469a3fd476b3c 100644 --- a/sycl/plugins/level_zero/pi_level_zero.hpp +++ b/sycl/plugins/level_zero/pi_level_zero.hpp @@ -940,11 +940,14 @@ struct _pi_buffer final : _pi_mem { } SubBuffer; }; +struct _pi_image; +using pi_image = _pi_image *; + // TODO: add proper support for images on context with multiple devices. struct _pi_image final : _pi_mem { // Image constructor - _pi_image(pi_context Ctx, ze_image_handle_t Image) - : _pi_mem(Ctx), ZeImage{Image} {} + _pi_image(pi_context Ctx, ze_image_handle_t Image, bool OwnNativeHandle) + : _pi_mem(Ctx), ZeImage{Image}, OwnZeMemHandle{OwnNativeHandle} {} virtual pi_result getZeHandle(char *&ZeHandle, access_mode_t, pi_device = nullptr) override { @@ -966,6 +969,8 @@ struct _pi_image final : _pi_mem { // Level Zero image handle. ze_image_handle_t ZeImage; + + bool OwnZeMemHandle; }; struct _pi_ze_event_list_t { diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 2c44f0cfe9eb3..24cfc9654539d 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -958,6 +958,21 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } +pi_result piextImgCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + bool ownNativeHandle, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + pi_mem *Img) { + (void)context; + (void)ownNativeHandle; + (void)ImageFormat; + (void)ImageDesc; + assert(Img != nullptr); + *Img = reinterpret_cast(nativeHandle); + return PI_SUCCESS; +} + pi_result piclProgramCreateWithSource(pi_context context, pi_uint32 count, const char **strings, const size_t *lengths, diff --git a/sycl/source/detail/image_impl.cpp b/sycl/source/detail/image_impl.cpp index 7ff987686496d..5be2ad656a8b2 100644 --- a/sycl/source/detail/image_impl.cpp +++ b/sycl/source/detail/image_impl.cpp @@ -297,6 +297,24 @@ image_impl::image_impl(cl_mem MemObject, const context &SyclContext, } } +image_impl::image_impl(pi_native_handle MemObject, const context &SyclContext, + event AvailableEvent, + std::unique_ptr Allocator, + uint8_t Dimensions, image_channel_order Order, + image_channel_type Type, bool OwnNativeHandle, + range<3> Range3WithOnes) + : BaseT(MemObject, SyclContext, OwnNativeHandle, std::move(AvailableEvent), + std::move(Allocator), detail::convertChannelOrder(Order), + detail::convertChannelType(Type), Range3WithOnes, Dimensions, + getImageElementSize(getImageNumberChannels(Order), Type)), + MDimensions(Dimensions), MRange(Range3WithOnes) { + MOrder = Order; + MType = Type; + MNumChannels = getImageNumberChannels(MOrder); + MElementSize = getImageElementSize(MNumChannels, Type); + setPitches(); // sets MRowPitch, MSlice and BaseT::MSizeInBytes +} + void *image_impl::allocateMem(ContextImplPtr Context, bool InitFromUserData, void *HostPtr, RT::PiEvent &OutEventToWait) { bool HostPtrReadOnly = false; diff --git a/sycl/source/detail/image_impl.hpp b/sycl/source/detail/image_impl.hpp index 9474fae5d895e..9a0f743272610 100644 --- a/sycl/source/detail/image_impl.hpp +++ b/sycl/source/detail/image_impl.hpp @@ -167,6 +167,12 @@ class __SYCL_EXPORT image_impl final : public SYCLMemObjT { std::unique_ptr Allocator, uint8_t Dimensions); + image_impl(pi_native_handle MemObject, const context &SyclContext, + event AvailableEvent, + std::unique_ptr Allocator, uint8_t Dimensions, + image_channel_order Order, image_channel_type Type, + bool OwnNativeHandle, range<3> Range3WithOnes); + // Return a range object representing the size of the image in terms of the // number of elements in each dimension as passed to the constructor range<3> get_range() const { return MRange; } diff --git a/sycl/source/detail/sycl_mem_obj_t.cpp b/sycl/source/detail/sycl_mem_obj_t.cpp index d64574bb6acac..57002366cbcbd 100644 --- a/sycl/source/detail/sycl_mem_obj_t.cpp +++ b/sycl/source/detail/sycl_mem_obj_t.cpp @@ -61,6 +61,65 @@ SYCLMemObjT::SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext, Plugin.call(MInteropMemObject); } +RT::PiMemObjectType getImageType(int Dimensions) { + if (Dimensions == 1) + return PI_MEM_TYPE_IMAGE1D; + if (Dimensions == 2) + return PI_MEM_TYPE_IMAGE2D; + return PI_MEM_TYPE_IMAGE3D; +} + +SYCLMemObjT::SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext, + bool OwnNativeHandle, event AvailableEvent, + std::unique_ptr Allocator, + RT::PiMemImageChannelOrder Order, + RT::PiMemImageChannelType Type, + range<3> Range3WithOnes, unsigned Dimensions, + size_t ElementSize) + : MAllocator(std::move(Allocator)), MProps(), + MInteropEvent(detail::getSyclObjImpl(std::move(AvailableEvent))), + MInteropContext(detail::getSyclObjImpl(SyclContext)), + MOpenCLInterop(true), MHostPtrReadOnly(false), MNeedWriteBack(true), + MUserPtr(nullptr), MShadowCopy(nullptr), MUploadDataFunctor(nullptr), + MSharedPtrStorage(nullptr), MHostPtrProvided(true) { + if (MInteropContext->is_host()) + throw sycl::invalid_parameter_error( + "Creation of interoperability memory object using host context is " + "not allowed", + PI_ERROR_INVALID_CONTEXT); + + RT::PiContext Context = nullptr; + const plugin &Plugin = getPlugin(); + + RT::PiMemImageFormat Format{Order, Type}; + RT::PiMemImageDesc Desc; + Desc.image_type = getImageType(Dimensions); + Desc.image_width = Range3WithOnes[0]; + Desc.image_height = Range3WithOnes[1]; + Desc.image_depth = Range3WithOnes[2]; + Desc.image_array_size = 0; + Desc.image_row_pitch = ElementSize * Desc.image_width; + Desc.image_slice_pitch = Desc.image_row_pitch * Desc.image_height; + Desc.num_mip_levels = 0; + Desc.num_samples = 0; + Desc.buffer = nullptr; + + Plugin.call( + MemObject, MInteropContext->getHandleRef(), OwnNativeHandle, &Format, + &Desc, &MInteropMemObject); + + Plugin.call(MInteropMemObject, PI_MEM_CONTEXT, + sizeof(Context), &Context, nullptr); + + if (MInteropContext->getHandleRef() != Context) + throw sycl::invalid_parameter_error( + "Input context must be the same as the context of cl_mem", + PI_ERROR_INVALID_CONTEXT); + + if (Plugin.getBackend() == backend::opencl) + Plugin.call(MInteropMemObject); +} + void SYCLMemObjT::releaseMem(ContextImplPtr Context, void *MemAllocation) { void *Ptr = getUserPtr(); return MemoryManager::releaseMemObj(Context, this, MemAllocation, Ptr); diff --git a/sycl/source/detail/sycl_mem_obj_t.hpp b/sycl/source/detail/sycl_mem_obj_t.hpp index dfd01b88c5a5a..5c0a6beb08994 100644 --- a/sycl/source/detail/sycl_mem_obj_t.hpp +++ b/sycl/source/detail/sycl_mem_obj_t.hpp @@ -79,6 +79,12 @@ class __SYCL_EXPORT SYCLMemObjT : public SYCLMemObjI { bool OwmNativeHandle, event AvailableEvent, std::unique_ptr Allocator); + SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext, + bool OwnNativeHandle, event AvailableEvent, + std::unique_ptr Allocator, + RT::PiMemImageChannelOrder Order, RT::PiMemImageChannelType Type, + range<3> Range3WithOnes, unsigned Dimensions, size_t ElementSize); + virtual ~SYCLMemObjT() = default; const plugin &getPlugin() const; diff --git a/sycl/source/image.cpp b/sycl/source/image.cpp index 781ff97e47e57..b98f29595125d 100644 --- a/sycl/source/image.cpp +++ b/sycl/source/image.cpp @@ -89,6 +89,17 @@ image_plain::image_plain(cl_mem ClMemObject, const context &SyclContext, } #endif +image_plain::image_plain(pi_native_handle MemObject, const context &SyclContext, + event AvailableEvent, + std::unique_ptr Allocator, + uint8_t Dimensions, image_channel_order Order, + image_channel_type Type, bool OwnNativeHandle, + range<3> Range3WithOnes) { + impl = std::make_shared( + MemObject, SyclContext, AvailableEvent, std::move(Allocator), Dimensions, + Order, Type, OwnNativeHandle, Range3WithOnes); +} + #define __SYCL_PARAM_TRAITS_SPEC(param_type) \ template <> \ __SYCL_EXPORT bool image_plain::has_property() const noexcept { \ diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index c2ac5e6863b8b..4730978ae1355 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -476,6 +476,15 @@ mock_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } +inline pi_result mock_piextImgCreateWithNativeHandle( + pi_native_handle NativeHandle, pi_context Context, bool OwnNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *RetImage) { + *RetImage = reinterpret_cast(NativeHandle); + retainDummyHandle(*RetImage); + return PI_SUCCESS; +} + // // Program // From 7bb157b74e76d665e0088c988d5e33f1482dbcde Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 9 Mar 2023 20:11:35 -0800 Subject: [PATCH 02/19] Linux ABI Symbols --- sycl/test/abi/sycl_symbols_linux.dump | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index 21f2d70806e09..c78126a8d60e1 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3,7 +3,7 @@ # DO NOT EDIT IT MANUALLY. Refer to sycl/doc/developer/ABIPolicyGuide.md for more info. ################################################################################ -# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so +# RUN: env LLVM_BIN_PATH=%llvm_build_bin_dir %python %sycl_tools_src_dir/abi_check.py --mode check_symbols --reference %s %sycl_libs_dir/libsycl.so # REQUIRES: linux # UNSUPPORTED: libcxx @@ -3722,7 +3722,9 @@ _ZN4sycl3_V16detail10image_impl11allocateMemESt10shared_ptrINS1_12context_implEE _ZN4sycl3_V16detail10image_impl14checkImageDescERK14_pi_image_descSt10shared_ptrINS1_12context_implEEPv _ZN4sycl3_V16detail10image_impl16checkImageFormatERK16_pi_image_formatSt10shared_ptrINS1_12context_implEE _ZN4sycl3_V16detail10image_implC1EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh +_ZN4sycl3_V16detail10image_implC1EmRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EEhNS0_19image_channel_orderENS0_18image_channel_typeEbNS0_5rangeILi3EEE _ZN4sycl3_V16detail10image_implC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh +_ZN4sycl3_V16detail10image_implC2EmRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EEhNS0_19image_channel_orderENS0_18image_channel_typeEbNS0_5rangeILi3EEE _ZN4sycl3_V16detail10make_eventEmRKNS0_7contextENS0_7backendE _ZN4sycl3_V16detail10make_eventEmRKNS0_7contextEbNS0_7backendE _ZN4sycl3_V16detail10make_queueEmRKNS0_7contextEPKNS0_6deviceEbRKSt8functionIFvNS0_14exception_listEEENS0_7backendE @@ -3733,8 +3735,10 @@ _ZN4sycl3_V16detail11SYCLMemObjT16updateHostMemoryEPv _ZN4sycl3_V16detail11SYCLMemObjT16updateHostMemoryEv _ZN4sycl3_V16detail11SYCLMemObjT20getBufSizeForContextERKSt10shared_ptrINS1_12context_implEEm _ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE +_ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE23_pi_image_channel_order22_pi_image_channel_typeNS0_5rangeILi3EEEjm _ZN4sycl3_V16detail11SYCLMemObjTC1EmRKNS0_7contextEmNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE _ZN4sycl3_V16detail11SYCLMemObjTC2EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE +_ZN4sycl3_V16detail11SYCLMemObjTC2EmRKNS0_7contextEbNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE23_pi_image_channel_order22_pi_image_channel_typeNS0_5rangeILi3EEEjm _ZN4sycl3_V16detail11SYCLMemObjTC2EmRKNS0_7contextEmNS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EE _ZN4sycl3_V16detail11buffer_impl11allocateMemESt10shared_ptrINS1_12context_implEEbPvRP9_pi_event _ZN4sycl3_V16detail11buffer_impl22destructorNotificationEPv @@ -3750,6 +3754,7 @@ _ZN4sycl3_V16detail11image_plainC1EPvNS0_19image_channel_orderENS0_18image_chann _ZN4sycl3_V16detail11image_plainC1EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EEhRKNS0_13property_listE _ZN4sycl3_V16detail11image_plainC1ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNSA_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISI_EEhRKNS0_13property_listEb _ZN4sycl3_V16detail11image_plainC1ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISF_EEhRKNS0_13property_listEb +_ZN4sycl3_V16detail11image_plainC1EmRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EEhNS0_19image_channel_orderENS0_18image_channel_typeEbNS0_5rangeILi3EEE _ZN4sycl3_V16detail11image_plainC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNS5_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISD_EEhRKNS0_13property_listE _ZN4sycl3_V16detail11image_plainC2ENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEhRKNS0_13property_listE _ZN4sycl3_V16detail11image_plainC2EP7_cl_memRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISA_EEh @@ -3758,6 +3763,7 @@ _ZN4sycl3_V16detail11image_plainC2EPvNS0_19image_channel_orderENS0_18image_chann _ZN4sycl3_V16detail11image_plainC2EPvNS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISB_EEhRKNS0_13property_listE _ZN4sycl3_V16detail11image_plainC2ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEERKNSA_ILi2EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISI_EEhRKNS0_13property_listEb _ZN4sycl3_V16detail11image_plainC2ERKSt10shared_ptrIKvENS0_19image_channel_orderENS0_18image_channel_typeERKNS0_5rangeILi3EEESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteISF_EEhRKNS0_13property_listEb +_ZN4sycl3_V16detail11image_plainC2EmRKNS0_7contextENS0_5eventESt10unique_ptrINS1_19SYCLMemObjAllocatorESt14default_deleteIS8_EEhNS0_19image_channel_orderENS0_18image_channel_typeEbNS0_5rangeILi3EEE _ZN4sycl3_V16detail11make_deviceEmNS0_7backendE _ZN4sycl3_V16detail11make_kernelERKNS0_7contextERKNS0_13kernel_bundleILNS0_12bundle_stateE2EEEmbNS0_7backendE _ZN4sycl3_V16detail11make_kernelEmRKNS0_7contextENS0_7backendE @@ -4454,7 +4460,6 @@ _ZNK4sycl3_V17context8get_infoINS0_4info7context32atomic_memory_scope_capabiliti _ZNK4sycl3_V17context8get_infoINS0_4info7context7devicesEEENS0_6detail20is_context_info_descIT_E11return_typeEv _ZNK4sycl3_V17context8get_infoINS0_4info7context8platformEEENS0_6detail20is_context_info_descIT_E11return_typeEv _ZNK4sycl3_V17context9getNativeEv -_ZNK4sycl3_V17handler17getContextImplPtrEv _ZNK4sycl3_V17handler27isStateExplicitKernelBundleEv _ZNK4sycl3_V17handler30getOrInsertHandlerKernelBundleEb _ZNK4sycl3_V17sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v From eb5a8a66b7b72d2bf5b55d2ec8589b041ac53116 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 9 Mar 2023 21:42:40 -0800 Subject: [PATCH 03/19] linux symbols revisited --- sycl/test/abi/sycl_symbols_linux.dump | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index c78126a8d60e1..9d9c10244b82d 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -4460,6 +4460,7 @@ _ZNK4sycl3_V17context8get_infoINS0_4info7context32atomic_memory_scope_capabiliti _ZNK4sycl3_V17context8get_infoINS0_4info7context7devicesEEENS0_6detail20is_context_info_descIT_E11return_typeEv _ZNK4sycl3_V17context8get_infoINS0_4info7context8platformEEENS0_6detail20is_context_info_descIT_E11return_typeEv _ZNK4sycl3_V17context9getNativeEv +_ZNK4sycl3_V17handler17getContextImplPtrEv _ZNK4sycl3_V17handler27isStateExplicitKernelBundleEv _ZNK4sycl3_V17handler30getOrInsertHandlerKernelBundleEb _ZNK4sycl3_V17sampler12get_propertyINS0_3ext6oneapi4cuda8property7context19use_primary_contextEEET_v From f80cb7ab6eec285b0bf27ad3b457099f4fe84b6e Mon Sep 17 00:00:00 2001 From: "Perkins, Chris" Date: Thu, 9 Mar 2023 21:53:29 -0800 Subject: [PATCH 04/19] windows symbols update. --- sycl/test/abi/sycl_symbols_windows.dump | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index edb7e962bd2bb..092f09c2c7235 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -376,6 +376,7 @@ ??0SYCLMemObjT@detail@_V1@sycl@@QEAA@PEAU_cl_mem@@AEBVcontext@23@Vevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@@Z ??0SYCLMemObjT@detail@_V1@sycl@@QEAA@_KAEBVcontext@23@_KVevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@@Z ??0SYCLMemObjT@detail@_V1@sycl@@QEAA@_KAEBVcontext@23@_NVevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@@Z +??0SYCLMemObjT@detail@_V1@sycl@@QEAA@_KAEBVcontext@23@_NVevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@W4_pi_image_channel_order@@W4_pi_image_channel_type@@V?$range@$02@23@I0@Z ??0SYCLMemObjT@detail@_V1@sycl@@QEAA@_KAEBVproperty_list@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@@Z ??0accelerator_selector@_V1@sycl@@QEAA@$$QEAV012@@Z ??0accelerator_selector@_V1@sycl@@QEAA@AEBV012@@Z @@ -483,6 +484,7 @@ ??0image_impl@detail@_V1@sycl@@QEAA@PEBXW4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z ??0image_impl@detail@_V1@sycl@@QEAA@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@AEBV?$range@$01@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z ??0image_impl@detail@_V1@sycl@@QEAA@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z +??0image_impl@detail@_V1@sycl@@QEAA@_KAEBVcontext@23@Vevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EW4image_channel_order@23@W4image_channel_type@23@_NV?$range@$02@23@@Z ??0image_plain@detail@_V1@sycl@@IEAA@AEBV?$shared_ptr@$$CBX@std@@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@AEBV?$range@$01@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@5@EAEBVproperty_list@23@_N@Z ??0image_plain@detail@_V1@sycl@@IEAA@AEBV?$shared_ptr@$$CBX@std@@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@5@EAEBVproperty_list@23@_N@Z ??0image_plain@detail@_V1@sycl@@IEAA@PEAU_cl_mem@@AEBVcontext@23@Vevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@E@Z @@ -491,6 +493,7 @@ ??0image_plain@detail@_V1@sycl@@IEAA@PEBXW4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z ??0image_plain@detail@_V1@sycl@@IEAA@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@AEBV?$range@$01@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z ??0image_plain@detail@_V1@sycl@@IEAA@W4image_channel_order@23@W4image_channel_type@23@AEBV?$range@$02@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EAEBVproperty_list@23@@Z +??0image_plain@detail@_V1@sycl@@IEAA@_KAEBVcontext@23@Vevent@23@V?$unique_ptr@VSYCLMemObjAllocator@detail@_V1@sycl@@U?$default_delete@VSYCLMemObjAllocator@detail@_V1@sycl@@@std@@@std@@EW4image_channel_order@23@W4image_channel_type@23@_NV?$range@$02@23@@Z ??0image_plain@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0image_plain@detail@_V1@sycl@@QEAA@AEBV0123@@Z ??0kernel@_V1@sycl@@AEAA@V?$shared_ptr@Vkernel_impl@detail@_V1@sycl@@@std@@@Z @@ -1015,8 +1018,8 @@ ?get_stream_mode@stream@_V1@sycl@@QEBA?AW4stream_manipulator@23@XZ ?get_wait_list@event@_V1@sycl@@QEAA?AV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@XZ ?get_width@stream@_V1@sycl@@QEBA_KXZ -?get_work_item_buffer_size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?get_work_item_buffer_size@stream@_V1@sycl@@QEBA_KXZ +?get_work_item_buffer_size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?gpu_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z ?handleHostData@SYCLMemObjT@detail@_V1@sycl@@QEAAXAEBV?$function@$$A6AXPEAX@Z@std@@_K_N@Z ?handleHostData@SYCLMemObjT@detail@_V1@sycl@@QEAAXAEBV?$shared_ptr@X@std@@_K_N@Z @@ -1184,8 +1187,8 @@ ?size@SYCLMemObjT@detail@_V1@sycl@@QEBA_KXZ ?size@exception_list@_V1@sycl@@QEBA_KXZ ?size@image_impl@detail@_V1@sycl@@QEBA_KXZ -?size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?size@stream@_V1@sycl@@QEBA_KXZ +?size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?split_string@detail@_V1@sycl@@YA?AV?$vector@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@V?$allocator@V?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@@2@@std@@AEBV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@5@D@Z ?start@HostProfilingInfo@detail@_V1@sycl@@QEAAXXZ ?start_fusion@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEAAXXZ From 20a0dead8c0ccbd314c35c7e2d768c2be10c6f25 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 10 Mar 2023 09:04:02 -0800 Subject: [PATCH 05/19] comments and doc update --- .../sycl_ext_oneapi_backend_level_zero.md | 27 +++++++++++++++++-- sycl/include/sycl/interop_handle.hpp | 4 +-- 2 files changed, 26 insertions(+), 5 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index c4100b6b39f51..b5f2593a4683d 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -425,10 +425,33 @@ make_image( This API is available starting with revision 4 of this specification. -Construct a SYCL image instance from a ze_image_handle_t. The input SYCL context Context must be associated with a single device, matching the device used at the prior allocation. +Construct a SYCL image instance from a ze_image_handle_t. + +Because LevelZero has no way of getting image information from an image, it must be provided. The backend_input_t is a struct type like so: +``` C++ +struct type { + ze_image_handle_t ZeImageHandle; + sycl::image_channel_order ChanOrder; + sycl::image_channel_type ChanType; + sycl::range Range; + ext::oneapi::level_zero::ownership Ownership{ + ext::oneapi::level_zero::ownership::transfer}; + }; +``` + +Example Usage +``` C++ +sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder, ChanType, ImgRange_2D, sycl::ext::oneapi::level_zero::ownership::transfer }; + +auto Image_2D = sycl::make_image(ImageInteropInput, Context); +``` + + The input SYCL context Context must be associated with a single device, matching the device used at the prior allocation. The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The Ownership input structure member 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. If the behavior is "transfer" then the runtime is going to free the input Level-Zero memory allocation. -Synchronization rules for a buffer that is created with this API are described in Section 4.5 +Synchronization rules with this API are described in Section 4.5 + + diff --git a/sycl/include/sycl/interop_handle.hpp b/sycl/include/sycl/interop_handle.hpp index b045b8b766597..5dad481ead77a 100644 --- a/sycl/include/sycl/interop_handle.hpp +++ b/sycl/include/sycl/interop_handle.hpp @@ -78,9 +78,7 @@ class interop_handle { /// accessor), the exception `sycl::invalid_object` is thrown /// asynchronously. template //, - // typename PropertyListT = ext::oneapi::accessor_property_list<>> + access::mode Mode, access::target Target, access::placeholder IsPlh> backend_return_t> get_native_mem( const detail::image_accessor Date: Fri, 10 Mar 2023 10:02:54 -0800 Subject: [PATCH 06/19] OCL and L0 ABI symbols --- sycl/test/abi/pi_level_zero_symbol_check.dump | 9 +++++---- sycl/test/abi/pi_opencl_symbol_check.dump | 9 +++++---- 2 files changed, 10 insertions(+), 8 deletions(-) diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index fbefe601f3675..151d2d04382e9 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -71,7 +71,6 @@ piProgramLink piProgramRelease piProgramRetain piQueueCreate -piextQueueCreate piQueueFinish piQueueFlush piQueueGetInfo @@ -92,6 +91,7 @@ piextDeviceSelectBinary piextEventCreateWithNativeHandle piextEventGetNativeHandle piextGetDeviceFunctionPointer +piextImgCreateWithNativeHandle piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj @@ -105,16 +105,17 @@ piextPluginGetOpaqueData piextProgramCreateWithNativeHandle piextProgramGetNativeHandle piextProgramSetSpecializationConstant +piextQueueCreate piextQueueCreateWithNativeHandle piextQueueGetNativeHandle piextUSMDeviceAlloc +piextUSMEnqueueFill2D piextUSMEnqueueMemAdvise piextUSMEnqueueMemcpy +piextUSMEnqueueMemcpy2D piextUSMEnqueueMemset -piextUSMEnqueuePrefetch -piextUSMEnqueueFill2D piextUSMEnqueueMemset2D -piextUSMEnqueueMemcpy2D +piextUSMEnqueuePrefetch piextUSMFree piextUSMGetMemAllocInfo piextUSMHostAlloc diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index 7925dfcbc6b53..f1f8b2ee28ff8 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -28,7 +28,6 @@ piProgramCreate piProgramCreateWithBinary piProgramLink piQueueCreate -piextQueueCreate piQueueGetInfo piSamplerCreate piTearDown @@ -40,6 +39,7 @@ piextDeviceGetNativeHandle piextDeviceSelectBinary piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer +piextImgCreateWithNativeHandle piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj @@ -52,16 +52,17 @@ piextPlatformGetNativeHandle piextProgramCreateWithNativeHandle piextProgramGetNativeHandle piextProgramSetSpecializationConstant +piextQueueCreate piextQueueCreateWithNativeHandle piextQueueGetNativeHandle piextUSMDeviceAlloc +piextUSMEnqueueFill2D piextUSMEnqueueMemAdvise piextUSMEnqueueMemcpy +piextUSMEnqueueMemcpy2D piextUSMEnqueueMemset -piextUSMEnqueuePrefetch -piextUSMEnqueueFill2D piextUSMEnqueueMemset2D -piextUSMEnqueueMemcpy2D +piextUSMEnqueuePrefetch piextUSMFree piextUSMGetMemAllocInfo piextUSMHostAlloc From 6db6a5ddf205ab1e3d0a7eb2539e4927477a3c2b Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 14 Mar 2023 09:30:21 -0700 Subject: [PATCH 07/19] documentation update --- .../extensions/supported/sycl_ext_oneapi_backend_level_zero.md | 1 + 1 file changed, 1 insertion(+) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index b5f2593a4683d..380b5385aad49 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -438,6 +438,7 @@ struct type { ext::oneapi::level_zero::ownership::transfer}; }; ``` +where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. Example Usage ``` C++ From 0dba3708752f2035b82382c0dca9cbf53215ae93 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Tue, 14 Mar 2023 09:36:12 -0700 Subject: [PATCH 08/19] reviewer feedback --- .../supported/sycl_ext_oneapi_backend_level_zero.md | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index 380b5385aad49..05b46f0a4aa2c 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -438,7 +438,8 @@ struct type { ext::oneapi::level_zero::ownership::transfer}; }; ``` -where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. +where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, +with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. Example Usage ``` C++ @@ -447,7 +448,7 @@ sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder auto Image_2D = sycl::make_image(ImageInteropInput, Context); ``` - The input SYCL context Context must be associated with a single device, matching the device used at the prior allocation. + The input SYCL context Context must be associated with a single device, matching the device used to create the Level Zero image handle. The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The Ownership input structure member 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. If the behavior is "transfer" then the runtime is going to free the input Level-Zero memory allocation. Synchronization rules with this API are described in Section 4.5 From c04317dd44aa8f7795360129d8311b500f807606 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 15 Mar 2023 09:07:08 -0700 Subject: [PATCH 09/19] comment and silence unused args --- sycl/plugins/cuda/pi_cuda.cpp | 19 ++++++++++--------- sycl/plugins/hip/pi_hip.cpp | 2 ++ 2 files changed, 12 insertions(+), 9 deletions(-) diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index fe56b4bd3bda5..09059eb6623a9 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2437,17 +2437,18 @@ pi_result cuda_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// TODO: Implement this. /// NOTE: The created PI object takes ownership of the native handle. /// -/// \param[in] nativeHandle The native handle to create PI mem object from. -/// \param[in] context The PI context of the memory allocation. -/// \param[in] ownNativeHandle Indicates if we own the native memory handle or -/// it came from interop that asked to not transfer the ownership to SYCL RT. -/// \param[out] mem Set to the PI mem object created from native handle. +/// \param[in] pi_native_handle The native handle to create PI mem object from. +/// \param[in] pi_context The PI context of the memory allocation. +/// \param[in] ownNativeHandle Boolean indicates if we own the native memory +/// handle or it came from interop that asked to not transfer the ownership to +/// SYCL RT. \param[in] pi_image_format The format of the image. \param[in] +/// pi_image_desc The description information for the image. \param[out] pi_mem +/// Set to the PI mem object created from native handle. /// /// \return TBD -pi_result cuda_piextImgCreateWithNativeHandle( - pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, - const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, - pi_mem *Img) { +pi_result cuda_piextImgCreateWithNativeHandle(pi_native_handle, pi_context, + bool, const pi_image_format *, + const pi_image_desc *, pi_mem *) { sycl::detail::pi::die( "Creation of PI mem from native image handle not implemented"); return {}; diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index dd0e31a938f67..c429d24b36347 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -2390,6 +2390,8 @@ pi_result hip_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// \param[in] context The PI context of the memory allocation. /// \param[in] ownNativeHandle Indicates if we own the native memory handle or /// it came from interop that asked to not transfer the ownership to SYCL RT. +/// \param[in] ImageFormat The format of the image. +/// \param[in] ImageDesc The description information for the image. /// \param[out] mem Set to the PI mem object created from native handle. /// /// \return TBD From dab95741334befc43bd572e73c5a4ada54415d9c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 17 Mar 2023 15:02:50 -0700 Subject: [PATCH 10/19] reviewer feedback --- .../sycl_ext_oneapi_backend_level_zero.md | 28 +++++++++++++++---- sycl/include/sycl/detail/pi.def | 2 +- sycl/include/sycl/detail/pi.h | 4 +-- sycl/plugins/cuda/pi_cuda.cpp | 7 +++-- .../esimd_emulator/pi_esimd_emulator.cpp | 6 ++-- sycl/plugins/hip/pi_hip.cpp | 10 +++---- sycl/plugins/level_zero/pi_level_zero.cpp | 12 ++++---- sycl/plugins/opencl/pi_opencl.cpp | 12 ++++---- sycl/source/detail/sycl_mem_obj_t.cpp | 2 +- sycl/test/abi/pi_level_zero_symbol_check.dump | 2 +- sycl/test/abi/pi_opencl_symbol_check.dump | 2 +- sycl/unittests/helpers/PiMockPlugin.hpp | 2 +- 12 files changed, 53 insertions(+), 36 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index 05b46f0a4aa2c..b2bd77da50b45 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -420,7 +420,7 @@ The additional AvailableEvent argument must be a valid SYCL event. make_image( const backend_input_t> &, - const context &Context, event AvailableEvent = {}) + const context &Context) ``` This API is available starting with revision 4 of this specification. @@ -440,20 +440,38 @@ struct type { ``` where the Range should be ordered (width), (width, height), or (width, height, depth) for 1D, 2D and 3D images respectively, with those values matching the dimensions used in the `ze_image_desc` that was used to create the `ze_image_handle_t` initially. +Note that the range term ordering (width first, depth last) is true for SYCL 1.2.1 images that are supported here. But future classes like +sampled_image and unsampled_image might have a different ordering. Example Usage ``` C++ sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder, ChanType, ImgRange_2D, sycl::ext::oneapi::level_zero::ownership::transfer }; -auto Image_2D = sycl::make_image(ImageInteropInput, Context); +sycl::image<2> Image_2D = sycl::make_image(ImageInteropInput, Context); ``` The input SYCL context Context must be associated with a single device, matching the device used to create the Level Zero image handle. -The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. -The Ownership input structure member 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. If the behavior is "transfer" then the runtime is going to free the input Level-Zero memory allocation. -Synchronization rules with this API are described in Section 4.5 +The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The created SYCL image can only be accessed from kernels that are submitted to a queue using this same context. +The Ownership input structure member 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. If the behavior is "transfer" then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. If the behavior is "keep", then the memory will not be freed by the ~image destructor, and it is the responsibility of the caller to free the memory appropriately. + + + + +``` C++ +make_image( + const backend_input_t> &, + const context &Context, event AvailableEvent) +``` + +This API is available starting with revision 4 of this specification. + +Construct a SYCL image instance from a pointer to a Level Zero memory allocation. Please refer to make_image +description above for semantics and restrictions. +The additional AvailableEvent argument must be a valid SYCL event. The instance of the SYCL image class template being constructed must wait for the SYCL event parameter to signal that the memory native handle is ready to be used. + diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index 11cf51f86be30..1dc0aefc2f9c4 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -61,7 +61,7 @@ _PI_API(piMemRelease) _PI_API(piMemBufferPartition) _PI_API(piextMemGetNativeHandle) _PI_API(piextMemCreateWithNativeHandle) -_PI_API(piextImgCreateWithNativeHandle) +_PI_API(piextMemImgCreateWithNativeHandle) // Program _PI_API(piProgramCreate) _PI_API(piclProgramCreateWithSource) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index ee9c126944add..79d44ff38d39c 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -77,7 +77,7 @@ // 12.22 Add piGetDeviceAndHostTimer to query device wall-clock timestamp // 12.23 Added new piextEnqueueDeviceGlobalVariableWrite and // piextEnqueueDeviceGlobalVariableRead functions. -// 12.24 Added piextImgCreateWithNativeHandle for creating images from native +// 12.24 Added piextMemImgCreateWithNativeHandle for creating images from native // handles. #define _PI_H_VERSION_MAJOR 12 @@ -1259,7 +1259,7 @@ __SYCL_EXPORT pi_result piextMemCreateWithNativeHandle( /// \param ownNativeHandle Indicates if we own the native memory handle or it /// came from interop that asked to not transfer the ownership to SYCL RT. /// \param img is the PI img created from the native handle. -__SYCL_EXPORT pi_result piextImgCreateWithNativeHandle( +__SYCL_EXPORT pi_result piextMemImgCreateWithNativeHandle( pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, pi_mem *img); diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 09059eb6623a9..0fc101e61504e 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2446,9 +2446,10 @@ pi_result cuda_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// Set to the PI mem object created from native handle. /// /// \return TBD -pi_result cuda_piextImgCreateWithNativeHandle(pi_native_handle, pi_context, - bool, const pi_image_format *, - const pi_image_desc *, pi_mem *) { +pi_result cuda_piextMemImgCreateWithNativeHandle(pi_native_handle, pi_context, + bool, const pi_image_format *, + const pi_image_desc *, + pi_mem *) { sycl::detail::pi::die( "Creation of PI mem from native image handle not implemented"); return {}; diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 93ffc7c5c3ae2..177480e26942f 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1297,9 +1297,9 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, DIE_NO_IMPLEMENTATION; } -pi_result piextImgCreateWithNativeHandle(pi_native_handle, pi_context, bool, - const pi_image_format *, - const pi_image_desc *, pi_mem *) { +pi_result piextMemImgCreateWithNativeHandle(pi_native_handle, pi_context, bool, + const pi_image_format *, + const pi_image_desc *, pi_mem *) { DIE_NO_IMPLEMENTATION; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index c429d24b36347..2e642f98d57b9 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -2395,12 +2395,10 @@ pi_result hip_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// \param[out] mem Set to the PI mem object created from native handle. /// /// \return TBD -pi_result hip_piextImgCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_context context, - bool ownNativeHandle, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - pi_mem *mem) { +pi_result hip_piextMemImgCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *mem) { (void)nativeHandle; (void)context; (void)ownNativeHandle; diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index e4fda0fa9a915..be6ccbe70e898 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3387,12 +3387,12 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle NativeHandle, return PI_SUCCESS; } -pi_result piextImgCreateWithNativeHandle(pi_native_handle NativeHandle, - pi_context Context, - bool OwnNativeHandle, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - pi_mem *RetImage) { +pi_result piextMemImgCreateWithNativeHandle(pi_native_handle NativeHandle, + pi_context Context, + bool OwnNativeHandle, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + pi_mem *RetImage) { PI_ASSERT(RetImage, PI_ERROR_INVALID_VALUE); PI_ASSERT(NativeHandle, PI_ERROR_INVALID_VALUE); diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 24cfc9654539d..eb71cf9b25754 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -958,12 +958,12 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } -pi_result piextImgCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_context context, - bool ownNativeHandle, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - pi_mem *Img) { +pi_result piextMemImgCreateWithNativeHandle(pi_native_handle nativeHandle, + pi_context context, + bool ownNativeHandle, + const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + pi_mem *Img) { (void)context; (void)ownNativeHandle; (void)ImageFormat; diff --git a/sycl/source/detail/sycl_mem_obj_t.cpp b/sycl/source/detail/sycl_mem_obj_t.cpp index 57002366cbcbd..22ebb53dcd14b 100644 --- a/sycl/source/detail/sycl_mem_obj_t.cpp +++ b/sycl/source/detail/sycl_mem_obj_t.cpp @@ -104,7 +104,7 @@ SYCLMemObjT::SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext, Desc.num_samples = 0; Desc.buffer = nullptr; - Plugin.call( + Plugin.call( MemObject, MInteropContext->getHandleRef(), OwnNativeHandle, &Format, &Desc, &MInteropMemObject); diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index 151d2d04382e9..824ddeb459055 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -91,7 +91,7 @@ piextDeviceSelectBinary piextEventCreateWithNativeHandle piextEventGetNativeHandle piextGetDeviceFunctionPointer -piextImgCreateWithNativeHandle +piextMemImgCreateWithNativeHandle piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index f1f8b2ee28ff8..016284b4afc5d 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -39,7 +39,7 @@ piextDeviceGetNativeHandle piextDeviceSelectBinary piextEventCreateWithNativeHandle piextGetDeviceFunctionPointer -piextImgCreateWithNativeHandle +piextMemImgCreateWithNativeHandle piextKernelCreateWithNativeHandle piextKernelGetNativeHandle piextKernelSetArgMemObj diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 4730978ae1355..0dcfc9c751fb8 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -476,7 +476,7 @@ mock_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } -inline pi_result mock_piextImgCreateWithNativeHandle( +inline pi_result mock_piextMemImgCreateWithNativeHandle( pi_native_handle NativeHandle, pi_context Context, bool OwnNativeHandle, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, pi_mem *RetImage) { From 45c973ce92ac9b3e41e3b7914071df1e544686a7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 3 Apr 2023 16:40:19 -0700 Subject: [PATCH 11/19] e2e tests --- ...nterop-level-zero-image-get-native-mem.cpp | 106 +++++++++ .../interop-level-zero-image-ownership.cpp | 137 +++++++++++ .../Plugin/interop-level-zero-image.cpp | 220 ++++++++++++++++++ 3 files changed, 463 insertions(+) create mode 100644 sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp create mode 100644 sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp create mode 100644 sycl/test-e2e/Plugin/interop-level-zero-image.cpp diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp new file mode 100644 index 0000000000000..aedadcf77db39 --- /dev/null +++ b/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp @@ -0,0 +1,106 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s + +// we use the interop to get the native image handle and then use that to make a +// new image and enumerate the pixels. + +// CHECK: (0 0) -- { 0 0 0 0 } +// CHECK-NEXT: (1 0) -- { 1 1 1 1 } +// CHECK-NEXT: (2 0) -- { 2 2 2 2 } +// CHECK-NEXT: (3 0) -- { 3 3 3 3 } +// CHECK-NEXT: (0 1) -- { 4 4 4 4 } +// CHECK-NEXT: (1 1) -- { 5 5 5 5 } +// CHECK-NEXT: (2 1) -- { 6 6 6 6 } +// CHECK-NEXT: (3 1) -- { 7 7 7 7 } + +// clang++ -fsycl -o las.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader +// interop-level-zero-image-get-native-mem.cpp + +#include +#include +#include +using namespace sycl; + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + sycl::device D = + sycl::ext::oneapi::filter_selector("level_zero:gpu").select_device(); + + sycl::context Ctx{D}; + sycl::queue Q(Ctx, D); + auto ZeContext = sycl::get_native(Ctx); + auto ZeDevice = sycl::get_native(D); + + // ----------- IMAGE STUFF + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + constexpr long width = 4; + constexpr long height = 2; + constexpr long numPixels = width * height; + ChannelDataT *sourceData = + (ChannelDataT *)std::calloc(numPixels * 4, sizeof(ChannelDataT)); + // initialize data: [ (0 0 0 0) (1 1 1 1) ...] + for (size_t i = 0; i < numPixels; i++) { + for (size_t chan = 0; chan < 4; chan++) { + size_t idx = (i * 4) + chan; + sourceData[idx] = (ChannelDataT)i; + } + } + // 8 bits per channel, four per pixel. + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + + const sycl::range<2> ImgRange_2D(width, height); + { // closure + // 1 - Create simple image. + sycl::image<2> image_2D(sourceData, ChanOrder, ChanType, ImgRange_2D); + + // 2 - Grab it's image handle via the get_native_mem interop. + using nativeH = sycl::backend_return_t>; + sycl::buffer passBack(range<1>{1}); + + Q.submit([&](handler &cgh) { + auto image_acc = + image_2D.get_access(cgh); + auto passBackAcc = passBack.get_host_access(sycl::write_only); + cgh.host_task([=](const interop_handle &IH) { + // There is nothing with image handles in the L0 API except + // create and destroy. So let's do that. + auto ZeImageH = IH.get_native_mem(image_acc); + passBackAcc[0] = ZeImageH; + }); + }).wait(); + + // Now we have the ZeImageH, so let's make a new SYCL image from it. + auto passBackAcc = passBack.get_host_access(sycl::read_only); + nativeH ZeImageH = passBackAcc[0]; + sycl::backend_input_t> imageData{ + ZeImageH, ChanOrder, ChanType, ImgRange_2D, + sycl::ext::oneapi::level_zero::ownership::keep}; + sycl::image<2> NewImg = sycl::make_image(imageData, Ctx); + + // Then use that image to read and stream out the data. + Q.submit([&](handler &cgh) { + auto read_acc = NewImg.get_access(cgh); + sycl::stream out(2024, 400, cgh); + cgh.single_task([=]() { + for (unsigned y = 0; y < height; y++) { + for (unsigned x = 0; x < width; x++) { + auto location = sycl::int2{x, y}; + pixelT somePixel = read_acc.read(location); + out << "(" << x << " " << y << ") -- { " << somePixel[0] << " " + << somePixel[1] << " " << somePixel[2] << " " << somePixel[3] + << " }" << sycl::endl; + } + } + }); + }).wait(); + } // ~image + +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + return 0; +} \ No newline at end of file diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp new file mode 100644 index 0000000000000..329ac64a47dbd --- /dev/null +++ b/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp @@ -0,0 +1,137 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: env ZE_DEBUG=1 %GPU_RUN_PLACEHOLDER %t.out 2>&1 | FileCheck %s + +// This test verifies that ownership is working correctly. +// If ownership is ::transfer then the ~image destructor will end up calling +// zeImageDestroy +// CHECK: test ownership::transfer +// CHECK: ZE ---> zeImageDestroy + +// With ownership ::keep it is must be called manually. +// CHECK: test ownership::keep +// CHECK: zeImageDestroy MANUAL + +// No other calls should appear. +// CHECK-NOT: zeImageDestroy + +// clang++ -fsycl -o wfd.bin -I$SYCL_HOME/build/install/include/sycl -lze_loader +// interop-level-zero-image-ownership.cpp + +#include +#include +#include + +using namespace sycl; + +void test(sycl::ext::oneapi::level_zero::ownership Ownership) { + + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + + platform Plt{gpu_selector_v}; + + auto Devices = Plt.get_devices(); + + if (Devices.size() < 1) { + std::cout << "Devices not found" << std::endl; + return; + } + + device Device = Devices[0]; + context Context{Device}; + queue Queue{Context, Device}; + + // Get native Level Zero handles + auto ZeContext = get_native(Context); + auto ZeDevice = get_native(Device); + + // ----------- Image Fundamentals + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + constexpr uint32_t numChannels = 4; // L0 only supports RGBA at this time. + + constexpr uint32_t width = 8; + constexpr uint32_t height = 4; + constexpr uint32_t depth = 1; + + const sycl::range<2> ImgRange_2D(width, height); + + // ----------- Basic LevelZero Description + ze_image_format_type_t ZeImageFormatType = ZE_IMAGE_FORMAT_TYPE_UINT; + size_t ZeImageFormatTypeSize = 8; + ze_image_format_layout_t ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + ze_image_format_t ZeFormatDesc = { + ZeImageFormatLayout, ZeImageFormatType, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}; + + ze_image_desc_t ZeImageDesc_base; + ZeImageDesc_base.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + ZeImageDesc_base.pNext = nullptr; + ZeImageDesc_base.flags = ZE_IMAGE_FLAG_KERNEL_WRITE; + // ZeImageDesc_base.flags = 0; + ZeImageDesc_base.arraylevels = 0; + ZeImageDesc_base.miplevels = 0; + ZeImageDesc_base.format = ZeFormatDesc; + + { + // ------ 2D ------ + ze_image_desc_t ZeImageDesc_2D = ZeImageDesc_base; + ZeImageDesc_2D.type = ZE_IMAGE_TYPE_2D; + ZeImageDesc_2D.width = width; + ZeImageDesc_2D.height = height; + ZeImageDesc_2D.depth = 1; + + ze_image_handle_t ZeHImage_2D; + ze_result_t res = + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_2D, &ZeHImage_2D); + if (res != ZE_RESULT_SUCCESS) { + std::cout << "unable to create image " << res << std::endl; + return; + } + + { // closure + sycl::backend_input_t> ImageInteropInput_2D{ + ZeHImage_2D, ChanOrder, ChanType, ImgRange_2D, Ownership}; + auto Image_2D = sycl::make_image(ImageInteropInput_2D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_2D.get_access(cgh); + + cgh.parallel_for(ImgRange_2D, [=](sycl::item<2> Item) { + auto location = sycl::int2{Item[0], Item[1]}; + auto sum = Item[0] + Item[1]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + } // ~image + // if ownership was transfer, then the ZeHImage_2D was destroyed as part of + // the ~image destruction (or deferred) + + if (Ownership == sycl::ext::oneapi::level_zero::ownership::keep) { + zeImageDestroy(ZeHImage_2D); + std::cout << "zeImageDestroy MANUAL" << std::endl; + } + + } // closure +} + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + std::cout << "test ownership::transfer" << std::endl; + test(sycl::ext::oneapi::level_zero::ownership::transfer); + + std::cout << "test ownership::keep" << std::endl; + test(sycl::ext::oneapi::level_zero::ownership::keep); +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + std::cout << "chau" << std::endl; + return 0; +} \ No newline at end of file diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image.cpp new file mode 100644 index 0000000000000..4c392b1e03582 --- /dev/null +++ b/sycl/test-e2e/Plugin/interop-level-zero-image.cpp @@ -0,0 +1,220 @@ +// REQUIRES: level_zero, level_zero_dev_kit +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %level_zero_options %s -o %t.out +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +// This test verifies that make_image is working for 1D, 2D and 3D images. +// We instantiate an image with L0, set its body, then use a host accessor to +// verify that the pixels are set correctly. + +// clang++ -fsycl -o ilzi.bin -I$SYCL_HOME/build/install/include/sycl +// -lze_loader interop-level-zero-image.cpp + +#include +#include +#include + +using namespace sycl; + +int main() { +#ifdef SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO + constexpr auto BE = sycl::backend::ext_oneapi_level_zero; + + platform Plt{gpu_selector_v}; + + auto Devices = Plt.get_devices(); + + if (Devices.size() < 1) { + std::cout << "Devices not found" << std::endl; + return 0; + } + + device Device = Devices[0]; + context Context{Device}; + queue Queue{Context, Device}; + + // Get native Level Zero handles + auto ZeContext = get_native(Context); + auto ZeDevice = get_native(Device); + + // ----------- Image Fundamentals + using pixelT = sycl::uint4; // accessor + using ChannelDataT = std::uint8_t; // allocator + sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; + sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; + constexpr uint32_t numChannels = 4; // L0 only supports RGBA at this time. + + constexpr uint32_t width = 8; + constexpr uint32_t height = 4; + constexpr uint32_t depth = 2; + + const sycl::range<1> ImgRange_1D(width); + const sycl::range<2> ImgRange_2D(width, height); + const sycl::range<3> ImgRange_3D(width, height, depth); + + // ----------- Basic LevelZero Description + ze_image_format_type_t ZeImageFormatType = ZE_IMAGE_FORMAT_TYPE_UINT; + size_t ZeImageFormatTypeSize = 8; + ze_image_format_layout_t ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8; + ze_image_format_t ZeFormatDesc = { + ZeImageFormatLayout, ZeImageFormatType, + ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G, + ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A}; + + ze_image_desc_t ZeImageDesc_base; + ZeImageDesc_base.stype = ZE_STRUCTURE_TYPE_IMAGE_DESC; + ZeImageDesc_base.pNext = nullptr; + ZeImageDesc_base.flags = ZE_IMAGE_FLAG_KERNEL_WRITE; + // ZeImageDesc_base.flags = 0; // <-- for read only + ZeImageDesc_base.arraylevels = 0; + ZeImageDesc_base.miplevels = 0; + ZeImageDesc_base.format = ZeFormatDesc; + + // ------ 1D ------ + { + std::cout << "glorious 1D" << std::endl; + // 1D image + ze_image_desc_t ZeImageDesc_1D = ZeImageDesc_base; + ZeImageDesc_1D.type = ZE_IMAGE_TYPE_1D; + ZeImageDesc_1D.width = width; + ZeImageDesc_1D.height = 1; + ZeImageDesc_1D.depth = 1; + + ze_image_handle_t ZeHImage_1D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_1D, &ZeHImage_1D); + + { // closure + sycl::backend_input_t> ImageInteropInput_1D{ + ZeHImage_1D, ChanOrder, ChanType, ImgRange_1D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_1D = sycl::make_image(ImageInteropInput_1D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_1D.get_access(cgh); + + cgh.parallel_for(ImgRange_1D, [=](sycl::item<1> Item) { + int x = Item[0]; + const pixelT somePixel = {x, x, x, x}; + write_acc.write(x, somePixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_1D.get_access(); + for (int col = 0; col < width; col++) { + const pixelT somePixel = read_acc.read(col); + // const pixelT expectedPixel = {col,col,col,col}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == col && somePixel[1] == col && + somePixel[2] == col && somePixel[3] == col); + } + + } // ~image + } // closure + + { + // ------ 2D ------ + std::cout << "glorious 2D" << std::endl; + // 2D image + ze_image_desc_t ZeImageDesc_2D = ZeImageDesc_base; + ZeImageDesc_2D.type = ZE_IMAGE_TYPE_2D; + ZeImageDesc_2D.width = width; + ZeImageDesc_2D.height = height; + ZeImageDesc_2D.depth = 1; + + ze_image_handle_t ZeHImage_2D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_2D, &ZeHImage_2D); + + { // closure + sycl::backend_input_t> ImageInteropInput_2D{ + ZeHImage_2D, ChanOrder, ChanType, ImgRange_2D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_2D = sycl::make_image(ImageInteropInput_2D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_2D.get_access(cgh); + + cgh.parallel_for(ImgRange_2D, [=](sycl::item<2> Item) { + auto location = sycl::int2{Item[0], Item[1]}; + auto sum = Item[0] + Item[1]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_2D.get_access(); + for (int row = 0; row < height; row++) { + for (int col = 0; col < width; col++) { + auto location = sycl::int2{col, row}; + const pixelT somePixel = read_acc.read(location); + auto sum = col + row; + // const pixelT expectedPixel = {sum,sum,sum,sum}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == sum && somePixel[1] == sum && + somePixel[2] == sum && somePixel[3] == sum); + } + } + + } // ~image + } // closure + + { + // ------ 3D ------ + std::cout << "glorious 3D" << std::endl; + // 3D image + ze_image_desc_t ZeImageDesc_3D = ZeImageDesc_base; + ZeImageDesc_3D.type = ZE_IMAGE_TYPE_3D; + ZeImageDesc_3D.width = width; + ZeImageDesc_3D.height = height; + ZeImageDesc_3D.depth = depth; + + ze_image_handle_t ZeHImage_3D; + zeImageCreate(ZeContext, ZeDevice, &ZeImageDesc_3D, &ZeHImage_3D); + + { // closure + sycl::backend_input_t> ImageInteropInput_3D{ + ZeHImage_3D, ChanOrder, ChanType, ImgRange_3D, + sycl::ext::oneapi::level_zero::ownership::keep}; + auto Image_3D = sycl::make_image(ImageInteropInput_3D, Context); + + Queue.submit([&](sycl::handler &cgh) { + auto write_acc = + Image_3D.get_access(cgh); + + cgh.parallel_for(ImgRange_3D, [=](sycl::item<3> Item) { + auto location = sycl::int4{Item[0], Item[1], Item[2], 0}; + auto sum = Item[0] + Item[1] + Item[2]; + const pixelT somepixel = {sum, sum, sum, sum}; + write_acc.write(location, somepixel); + }); + }); + Queue.wait_and_throw(); + + // now check with host accessor. + auto read_acc = Image_3D.get_access(); + for (int row = 0; row < height; row++) { + for (int col = 0; col < width; col++) { + for (int z = 0; z < depth; z++) { + auto location = sycl::int4{col, row, z, 0}; + const pixelT somePixel = read_acc.read(location); + auto sum = col + row + z; + // const pixelT expectedPixel = {sum,sum,sum,sum}; + // assert(somePixel == expectedPixel); + assert(somePixel[0] == sum && somePixel[1] == sum && + somePixel[2] == sum && somePixel[3] == sum); + } + } + } + + } // ~image + } // closure + +#else + std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; +#endif + return 0; +} \ No newline at end of file From 401e476a3006e82982bb086203aac4f90140e9a7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 3 Apr 2023 19:37:48 -0700 Subject: [PATCH 12/19] reviewer doc feedback --- .../supported/sycl_ext_oneapi_backend_level_zero.md | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index b2bd77da50b45..22ecb2dbc29c5 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -417,7 +417,7 @@ The additional AvailableEvent argument must be a valid SYCL event. ``` C++ -make_image( +make_image( const backend_input_t> &, const context &Context) @@ -452,7 +452,7 @@ sycl::image<2> Image_2D = sycl::make_image(ImageInteropInput, Context); The input SYCL context Context must be associated with a single device, matching the device used to create the Level Zero image handle. The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The created SYCL image can only be accessed from kernels that are submitted to a queue using this same context. -The Ownership input structure member 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. If the behavior is "transfer" then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. If the behavior is "keep", then the memory will not be freed by the ~image destructor, and it is the responsibility of the caller to free the memory appropriately. +The Ownership input structure member 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. If the behavior is "transfer" then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. When using "transfer" the ~image destructor may not need to block. If the behavior is "keep", then the memory will not be freed by the ~image destructor, and the ~image destructor blocks until all work in the queues on the image have been completed. When using "keep" it is the responsibility of the caller to free the memory appropriately. From e5879179c166b6674a23c2dd1e28f2d70edf268e Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 6 Apr 2023 17:44:22 -0700 Subject: [PATCH 13/19] reviewer feedback Signed-off-by: Chris Perkins --- .../sycl_ext_oneapi_backend_level_zero.md | 30 +++++++++++-------- sycl/include/sycl/detail/pi.def | 2 +- sycl/include/sycl/detail/pi.h | 13 +++++--- sycl/plugins/cuda/pi_cuda.cpp | 9 +++--- .../esimd_emulator/pi_esimd_emulator.cpp | 6 ++-- sycl/plugins/hip/pi_hip.cpp | 2 +- sycl/plugins/level_zero/pi_level_zero.cpp | 20 ++++++------- sycl/plugins/opencl/pi_opencl.cpp | 10 +++---- sycl/source/detail/sycl_mem_obj_t.cpp | 2 +- sycl/test/abi/pi_level_zero_symbol_check.dump | 2 +- sycl/test/abi/pi_opencl_symbol_check.dump | 2 +- sycl/unittests/helpers/PiMockPlugin.hpp | 2 +- 12 files changed, 54 insertions(+), 46 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index f50564dd329f7..fbca56a863de2 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -263,7 +263,7 @@ struct { ze_image_handle_t ZeImageHandle; sycl::image_channel_order ChanOrder; sycl::image_channel_type ChanType; - range Range; + sycl::range Range; ext::oneapi::level_zero::ownership Ownership{ ext::oneapi::level_zero::ownership::transfer}; } @@ -461,13 +461,13 @@ The additional AvailableEvent argument must be a valid SYCL event. ``` C++ -make_image( - const backend_input_t> &, - const context &Context) +template +image make_image( + const backend_input_t> &backendObject, + const context &targetContext); ``` -This API is available starting with revision 4 of this specification. +This API is available starting with revision 5 of this specification. Construct a SYCL image instance from a ze_image_handle_t. @@ -489,9 +489,15 @@ sampled_image and unsampled_image might have a different ordering. Example Usage ``` C++ -sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder, ChanType, ImgRange_2D, sycl::ext::oneapi::level_zero::ownership::transfer }; +sycl::backend_input_t> ImageInteropInput{ + ZeHImage, + ChanOrder, + ChanType, + ImgRange_2D, + sycl::ext::oneapi::level_zero::ownership::transfer }; -sycl::image<2> Image_2D = sycl::make_image(ImageInteropInput, Context); +sycl::image<2> Image_2D + = sycl::make_image(ImageInteropInput, Context); ``` The input SYCL context Context must be associated with a single device, matching the device used to create the Level Zero image handle. @@ -504,10 +510,10 @@ The Ownership input structure member specifies if the SYCL runtime ``` C++ -make_image( - const backend_input_t> &, - const context &Context, event AvailableEvent) +template +image make_image( + const backend_input_t> &backendObject, + const context &targetContext, event availableEvent); ``` This API is available starting with revision 4 of this specification. diff --git a/sycl/include/sycl/detail/pi.def b/sycl/include/sycl/detail/pi.def index 9a9e2d8d6c430..fd5bc5a844cef 100644 --- a/sycl/include/sycl/detail/pi.def +++ b/sycl/include/sycl/detail/pi.def @@ -61,7 +61,7 @@ _PI_API(piMemRelease) _PI_API(piMemBufferPartition) _PI_API(piextMemGetNativeHandle) _PI_API(piextMemCreateWithNativeHandle) -_PI_API(piextMemImgCreateWithNativeHandle) +_PI_API(piextMemImageCreateWithNativeHandle) // Program _PI_API(piProgramCreate) _PI_API(piclProgramCreateWithSource) diff --git a/sycl/include/sycl/detail/pi.h b/sycl/include/sycl/detail/pi.h index 7c7e5eb48c831..38d173137c82e 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -86,8 +86,8 @@ // 12.27 Added new queue create and get APIs for immediate commandlists // piextQueueCreate2, piextQueueCreateWithNativeHandle2, // piextQueueGetNativeHandle2 -// 12.28 Added piextMemImgCreateWithNativeHandle for creating images from native -// handles. +// 12.28 Added piextMemImageCreateWithNativeHandle for creating images from +// native handles. #define _PI_H_VERSION_MAJOR 12 #define _PI_H_VERSION_MINOR 28 @@ -1311,14 +1311,19 @@ __SYCL_EXPORT pi_result piextMemCreateWithNativeHandle( pi_mem *mem); /// Creates PI image object from a native handle. -/// NOTE: The created PI object takes ownership of the native handle. /// /// \param nativeHandle is the native handle to create PI image from. /// \param context The PI context of the memory allocation. /// \param ownNativeHandle Indicates if we own the native memory handle or it /// came from interop that asked to not transfer the ownership to SYCL RT. +/// \param ImageFormat is the pi_image_format struct that +/// specifies the image channnel order and channel data type that +/// match what the nativeHandle uses +/// \param ImageDesc is the pi_image_desc struct that specifies +/// the image dimension, pitch, slice and other information about +/// the nativeHandle /// \param img is the PI img created from the native handle. -__SYCL_EXPORT pi_result piextMemImgCreateWithNativeHandle( +__SYCL_EXPORT pi_result piextMemImageCreateWithNativeHandle( pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, pi_mem *img); diff --git a/sycl/plugins/cuda/pi_cuda.cpp b/sycl/plugins/cuda/pi_cuda.cpp index 2c5a2c6440aa4..68b4fb8ef60b9 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2465,10 +2465,11 @@ pi_result cuda_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// Set to the PI mem object created from native handle. /// /// \return TBD -pi_result cuda_piextMemImgCreateWithNativeHandle(pi_native_handle, pi_context, - bool, const pi_image_format *, - const pi_image_desc *, - pi_mem *) { +pi_result cuda_piextMemImageCreateWithNativeHandle(pi_native_handle, pi_context, + bool, + const pi_image_format *, + const pi_image_desc *, + pi_mem *) { sycl::detail::pi::die( "Creation of PI mem from native image handle not implemented"); return {}; diff --git a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp index 478199fe791ef..ed0fd2fa72c9d 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1314,9 +1314,9 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, DIE_NO_IMPLEMENTATION; } -pi_result piextMemImgCreateWithNativeHandle(pi_native_handle, pi_context, bool, - const pi_image_format *, - const pi_image_desc *, pi_mem *) { +pi_result piextMemImageCreateWithNativeHandle(pi_native_handle, pi_context, + bool, const pi_image_format *, + const pi_image_desc *, pi_mem *) { DIE_NO_IMPLEMENTATION; } diff --git a/sycl/plugins/hip/pi_hip.cpp b/sycl/plugins/hip/pi_hip.cpp index 084f5427f05dc..311af5da41ab4 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -2464,7 +2464,7 @@ pi_result hip_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, /// \param[out] mem Set to the PI mem object created from native handle. /// /// \return TBD -pi_result hip_piextMemImgCreateWithNativeHandle( +pi_result hip_piextMemImageCreateWithNativeHandle( pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, pi_mem *mem) { diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 55f80c6704048..7ecd46c0af667 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3273,9 +3273,9 @@ pi_result piMemRelease(pi_mem Mem) { return PI_SUCCESS; } -pi_result PIToZeImageDesc(const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - ZeStruct &ZeImageDesc) { +static pi_result pi2zeImageDesc(const pi_image_format *ImageFormat, + const pi_image_desc *ImageDesc, + ZeStruct &ZeImageDesc) { ze_image_format_type_t ZeImageFormatType; size_t ZeImageFormatTypeSize; switch (ImageFormat->image_channel_data_type) { @@ -3415,7 +3415,7 @@ pi_result piMemImageCreate(pi_context Context, pi_mem_flags Flags, ZeStruct ZeImageDesc; pi_result DescriptionResult = - PIToZeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); + pi2zeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); if (DescriptionResult != PI_SUCCESS) return DescriptionResult; @@ -3561,12 +3561,10 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle NativeHandle, return PI_SUCCESS; } -pi_result piextMemImgCreateWithNativeHandle(pi_native_handle NativeHandle, - pi_context Context, - bool OwnNativeHandle, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - pi_mem *RetImage) { +pi_result piextMemImageCreateWithNativeHandle( + pi_native_handle NativeHandle, pi_context Context, bool OwnNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *RetImage) { PI_ASSERT(RetImage, PI_ERROR_INVALID_VALUE); PI_ASSERT(NativeHandle, PI_ERROR_INVALID_VALUE); @@ -3583,7 +3581,7 @@ pi_result piextMemImgCreateWithNativeHandle(pi_native_handle NativeHandle, #ifndef NDEBUG ZeStruct ZeImageDesc; pi_result DescriptionResult = - PIToZeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); + pi2zeImageDesc(ImageFormat, ImageDesc, ZeImageDesc); if (DescriptionResult != PI_SUCCESS) return DescriptionResult; diff --git a/sycl/plugins/opencl/pi_opencl.cpp b/sycl/plugins/opencl/pi_opencl.cpp index 223710dd21df8..48fd7dc5017db 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -1212,12 +1212,10 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } -pi_result piextMemImgCreateWithNativeHandle(pi_native_handle nativeHandle, - pi_context context, - bool ownNativeHandle, - const pi_image_format *ImageFormat, - const pi_image_desc *ImageDesc, - pi_mem *Img) { +pi_result piextMemImageCreateWithNativeHandle( + pi_native_handle nativeHandle, pi_context context, bool ownNativeHandle, + const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, + pi_mem *Img) { (void)context; (void)ownNativeHandle; (void)ImageFormat; diff --git a/sycl/source/detail/sycl_mem_obj_t.cpp b/sycl/source/detail/sycl_mem_obj_t.cpp index 22ebb53dcd14b..95022c60d72ce 100644 --- a/sycl/source/detail/sycl_mem_obj_t.cpp +++ b/sycl/source/detail/sycl_mem_obj_t.cpp @@ -104,7 +104,7 @@ SYCLMemObjT::SYCLMemObjT(pi_native_handle MemObject, const context &SyclContext, Desc.num_samples = 0; Desc.buffer = nullptr; - Plugin.call( + Plugin.call( MemObject, MInteropContext->getHandleRef(), OwnNativeHandle, &Format, &Desc, &MInteropMemObject); diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index 945dd38efa6a8..08292e4fbeb45 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -100,7 +100,7 @@ piextKernelSetArgPointer piextKernelSetArgSampler piextMemCreateWithNativeHandle piextMemGetNativeHandle -piextMemImgCreateWithNativeHandle +piextMemImageCreateWithNativeHandle piextPlatformCreateWithNativeHandle piextPlatformGetNativeHandle piextPluginGetOpaqueData diff --git a/sycl/test/abi/pi_opencl_symbol_check.dump b/sycl/test/abi/pi_opencl_symbol_check.dump index db3e187748c49..a367ec1afd33d 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -48,7 +48,7 @@ piextKernelSetArgPointer piextKernelSetArgSampler piextMemCreateWithNativeHandle piextMemGetNativeHandle -piextMemImgCreateWithNativeHandle +piextMemImageCreateWithNativeHandle piextPlatformCreateWithNativeHandle piextPlatformGetNativeHandle piextProgramCreateWithNativeHandle diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index e4523c1d75380..a2fd198c4055d 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -498,7 +498,7 @@ mock_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } -inline pi_result mock_piextMemImgCreateWithNativeHandle( +inline pi_result mock_piextMemImageCreateWithNativeHandle( pi_native_handle NativeHandle, pi_context Context, bool OwnNativeHandle, const pi_image_format *ImageFormat, const pi_image_desc *ImageDesc, pi_mem *RetImage) { From 34c38d1dc282f9db44c5bbb881a6fb50773c28e9 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Fri, 7 Apr 2023 10:23:49 -0700 Subject: [PATCH 14/19] moar reviewer feedback Signed-off-by: Chris Perkins --- .../sycl_ext_oneapi_backend_level_zero.md | 65 +++++++++++++------ sycl/plugins/level_zero/pi_level_zero.cpp | 3 +- 2 files changed, 48 insertions(+), 20 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index fbca56a863de2..3ea3d4432fa4b 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -322,7 +322,8 @@ an application to create a SYCL object that encapsulates a corresponding Level-Z ``` C++ -make_platform( +template +platform make_platform( const backend_input_t< backend::ext_oneapi_level_zero, platform> &) ``` @@ -332,7 +333,8 @@ make_platform( ``` C++ -make_device( +template +device make_device( const backend_input_t< backend::ext_oneapi_level_zero, device> &) ``` @@ -342,7 +344,8 @@ make_device( ``` C++ -make_context( +template +context make_context( const backend_input_t< backend::ext_oneapi_level_zero, context> &) ``` @@ -352,7 +355,8 @@ make_context( ``` C++ -make_queue( +template +queue make_queue( const backend_input_t< backend::ext_oneapi_level_zero, queue> &, const context &Context) @@ -369,7 +373,8 @@ the ```compute_index``` property which is built into the command queue or comman ``` C++ -make_event( +template +event make_event( const backend_input_t< backend::ext_oneapi_level_zero, event> &, const context &Context) @@ -380,11 +385,11 @@ make_event( ``` C++ -make_kernel_bundle( +template +kernel_bundle make_kernel_bundle( const backend_input_t< backend::ext_oneapi_level_zero, - kernel_bundle> &, + kernel_bundle> &, const context &Context) ``` @@ -406,7 +411,8 @@ interoperability kernel_bundle destructor is called. ``` C++ -make_kernel( +template +kernel make_kernel( const backend_input_t< backend::ext_oneapi_level_zero, kernel> &, const context &Context) @@ -428,8 +434,11 @@ Level-Zero kernel ``` C++ -make_buffer( - const backend_input_t>> +buffer make_buffer( + const backend_input_t> &, const context &Context) ``` @@ -444,8 +453,11 @@ Synchronization rules for a buffer that is created with this API are described i ``` C++ -make_buffer( - const backend_input_t>> +buffer make_buffer( + const backend_input_t> &, const context &Context, event AvailableEvent) ``` @@ -461,9 +473,11 @@ The additional AvailableEvent argument must be a valid SYCL event. ``` C++ -template +template image make_image( - const backend_input_t> &backendObject, + const backend_input_t> &backendObject, const context &targetContext); ``` @@ -489,6 +503,17 @@ sampled_image and unsampled_image might have a different ordering. Example Usage ``` C++ +ze_image_handle_t ZeHImage; +// ... user provided LevelZero ZeHImage image handle gotten somehow (possibly zeImageCreate) + +// the informational data that matches ZeHImage +sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; +sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; +constexpr uint32_t width = 4; +constexpr uint32_t height = 2; +sycl::range<2> ImgRange_2D(width, height); + +constexpr sycl::backend BE = sycl::backend::ext_oneapi_level_zero; sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder, @@ -500,7 +525,7 @@ sycl::image<2> Image_2D = sycl::make_image(ImageInteropInput, Context); ``` - The input SYCL context Context must be associated with a single device, matching the device used to create the Level Zero image handle. +The image can only be used on the single device where it was created. This limitation may be relaxed in the future. The Context argument must be a valid SYCL context encapsulating a Level-Zero context, and the Level-Zero image must have been created on the same context. The created SYCL image can only be accessed from kernels that are submitted to a queue using this same context. The Ownership input structure member 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. If the behavior is "transfer" then the SYCL runtime is going to free the input Level-Zero memory allocation, meaning the memory will be freed when the ~image destructor fires. When using "transfer" the ~image destructor may not need to block. If the behavior is "keep", then the memory will not be freed by the ~image destructor, and the ~image destructor blocks until all work in the queues on the image have been completed. When using "keep" it is the responsibility of the caller to free the memory appropriately. @@ -510,13 +535,15 @@ The Ownership input structure member specifies if the SYCL runtime ``` C++ -template +template image make_image( - const backend_input_t> &backendObject, + const backend_input_t> &backendObject, const context &targetContext, event availableEvent); ``` -This API is available starting with revision 4 of this specification. +This API is available starting with revision 5 of this specification. Construct a SYCL image instance from a pointer to a Level Zero memory allocation. Please refer to make_image description above for semantics and restrictions. diff --git a/sycl/plugins/level_zero/pi_level_zero.cpp b/sycl/plugins/level_zero/pi_level_zero.cpp index 7ecd46c0af667..31ef8e7df2b43 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3259,7 +3259,8 @@ pi_result piMemRelease(pi_mem Mem) { auto Image = static_cast(Mem); if (Image->OwnZeMemHandle) { PI_CALL(Mem->getZeHandle(ZeHandleImage, _pi_mem::write_only)); - auto ZeResult = ZE_CALL_NOCHECK(zeImageDestroy, (pi_cast(ZeHandleImage))); + auto ZeResult = ZE_CALL_NOCHECK( + zeImageDestroy, (pi_cast(ZeHandleImage))); // Gracefully handle the case that L0 was already unloaded. if (ZeResult && ZeResult != ZE_RESULT_ERROR_UNINITIALIZED) return mapError(ZeResult); From 423058a7b0a3e44dc073e46db569b1c0fcf48d7c Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Mon, 10 Apr 2023 08:44:43 -0700 Subject: [PATCH 15/19] more doc changes, reviewer feedback and spacing --- .../sycl_ext_oneapi_backend_level_zero.md | 38 +++++++++---------- 1 file changed, 18 insertions(+), 20 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index 3ea3d4432fa4b..cebbc714460d0 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -324,8 +324,7 @@ an application to create a SYCL object that encapsulates a corresponding Level-Z ``` C++ template platform make_platform( - const backend_input_t< - backend::ext_oneapi_level_zero, platform> &) + const backend_input_t &) ``` Constructs a SYCL platform instance from a Level-Zero ze_driver_handle_t. The SYCL execution environment contains a fixed number of platforms that are enumerated via sycl::platform::get_platforms(). Calling this function does not create a new platform. Rather it merely creates a sycl::platform object that is a copy of one of the platforms from that enumeration. @@ -335,8 +334,7 @@ platform make_platform( ``` C++ template device make_device( - const backend_input_t< - backend::ext_oneapi_level_zero, device> &) + const backend_input_t &) ``` Constructs a SYCL device instance from a Level-Zero ze_device_handle_t. The SYCL execution environment for the Level Zero backend contains a fixed number of devices that are enumerated via sycl::device::get_devices() and a fixed number of sub-devices that are enumerated via sycl::device::create_sub_devices(...). Calling this function does not create a new device. Rather it merely creates a sycl::device object that is a copy of one of the devices from those enumerations. @@ -346,8 +344,7 @@ device make_device( ``` C++ template context make_context( - const backend_input_t< - backend::ext_oneapi_level_zero, context> &) + const backend_input_t &) ``` Constructs a SYCL context instance from a Level-Zero ze_context_handle_t. The context is created against the devices passed in DeviceList structure member. 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 input structure member 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. @@ -357,8 +354,7 @@ context make_context( ``` C++ template queue make_queue( - const backend_input_t< - backend::ext_oneapi_level_zero, queue> &, + const backend_input_t &, const context &Context) ``` @@ -375,8 +371,7 @@ the ```compute_index``` property which is built into the command queue or comman ``` C++ template event make_event( - const backend_input_t< - backend::ext_oneapi_level_zero, event> &, + const backend_input_t &, const context &Context) ``` @@ -385,10 +380,10 @@ event make_event( ``` C++ +// State must be bundle_state::executable template kernel_bundle make_kernel_bundle( - const backend_input_t< - backend::ext_oneapi_level_zero, + const backend_input_t> &, const context &Context) ``` @@ -413,8 +408,7 @@ interoperability kernel_bundle destructor is called. ``` C++ template kernel make_kernel( - const backend_input_t< - backend::ext_oneapi_level_zero, kernel> &, + const backend_input_t &, const context &Context) ``` @@ -504,16 +498,20 @@ sampled_image and unsampled_image might have a different ordering. Example Usage ``` C++ ze_image_handle_t ZeHImage; -// ... user provided LevelZero ZeHImage image handle gotten somehow (possibly zeImageCreate) +// ... user provided LevelZero ZeHImage image +// handle gotten somehow (possibly zeImageCreate) // the informational data that matches ZeHImage -sycl::image_channel_order ChanOrder = sycl::image_channel_order::rgba; -sycl::image_channel_type ChanType = sycl::image_channel_type::unsigned_int8; -constexpr uint32_t width = 4; -constexpr uint32_t height = 2; +sycl::image_channel_order ChanOrder + = sycl::image_channel_order::rgba; +sycl::image_channel_type ChanType + = sycl::image_channel_type::unsigned_int8; +size_t width = 4; +size_t height = 2; sycl::range<2> ImgRange_2D(width, height); -constexpr sycl::backend BE = sycl::backend::ext_oneapi_level_zero; +constexpr sycl::backend BE + = sycl::backend::ext_oneapi_level_zero; sycl::backend_input_t> ImageInteropInput{ ZeHImage, ChanOrder, From 8b33f1b11584e633435a6db45709e68e332816cf Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Apr 2023 13:10:30 -0700 Subject: [PATCH 16/19] diet and excercise --- .../sycl_ext_oneapi_backend_level_zero.md | 14 ++++++++------ 1 file changed, 8 insertions(+), 6 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index cebbc714460d0..e22804959bd74 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -468,10 +468,11 @@ The additional AvailableEvent argument must be a valid SYCL event. ``` C++ template -image make_image( + typename AllocrT = sycl::image_allocator> +image make_image( const backend_input_t> &backendObject, + image> &backendObject, const context &targetContext); ``` @@ -534,10 +535,11 @@ The Ownership input structure member specifies if the SYCL runtime ``` C++ template -image make_image( + typename AllocrT = sycl::image_allocator> +image make_image( const backend_input_t> &backendObject, + image> &backendObject, const context &targetContext, event availableEvent); ``` From 60191f6df9060bf0a0db28813d18f998aee1120d Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Apr 2023 13:16:14 -0700 Subject: [PATCH 17/19] more reformat to elim scroll bars in github preview --- .../sycl_ext_oneapi_backend_level_zero.md | 14 ++++++++++---- 1 file changed, 10 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index e22804959bd74..bbb062ab12d1f 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -430,10 +430,13 @@ Level-Zero kernel ``` C++ template >> + typename AllocatorT = + buffer_allocator>> buffer make_buffer( const backend_input_t> &, + buffer> &, const context &Context) ``` @@ -449,10 +452,13 @@ Synchronization rules for a buffer that is created with this API are described i ``` C++ template >> + typename AllocatorT = + buffer_allocator>> buffer make_buffer( const backend_input_t> &, + buffer> &, const context &Context, event AvailableEvent) ``` From 614f85a0eb41ee68f3c45f9707fc43d84a31c2c7 Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Wed, 12 Apr 2023 13:29:42 -0700 Subject: [PATCH 18/19] removed unneeded specializations from doc --- .../sycl_ext_oneapi_backend_level_zero.md | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md index bbb062ab12d1f..297e1a53d2e53 100644 --- a/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md +++ b/sycl/doc/extensions/supported/sycl_ext_oneapi_backend_level_zero.md @@ -323,7 +323,7 @@ an application to create a SYCL object that encapsulates a corresponding Level-Z ``` C++ template -platform make_platform( +platform make_platform( const backend_input_t &) ``` @@ -333,7 +333,7 @@ platform make_platform( ``` C++ template -device make_device( +device make_device( const backend_input_t &) ``` @@ -343,7 +343,7 @@ device make_device( ``` C++ template -context make_context( +context make_context( const backend_input_t &) ``` @@ -353,7 +353,7 @@ context make_context( ``` C++ template -queue make_queue( +queue make_queue( const backend_input_t &, const context &Context) ``` @@ -370,7 +370,7 @@ the ```compute_index``` property which is built into the command queue or comman ``` C++ template -event make_event( +event make_event( const backend_input_t &, const context &Context) ``` @@ -382,7 +382,7 @@ event make_event( ``` C++ // State must be bundle_state::executable template -kernel_bundle make_kernel_bundle( +kernel_bundle make_kernel_bundle( const backend_input_t> &, const context &Context) @@ -407,7 +407,7 @@ interoperability kernel_bundle destructor is called. ``` C++ template -kernel make_kernel( +kernel make_kernel( const backend_input_t &, const context &Context) ``` From 50fea339a92b960bc3bc59ccfec3a8cc897cde9f Mon Sep 17 00:00:00 2001 From: Chris Perkins Date: Thu, 13 Apr 2023 14:51:43 -0700 Subject: [PATCH 19/19] add newline to end of tests --- .../test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp | 2 +- sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp | 2 +- sycl/test-e2e/Plugin/interop-level-zero-image.cpp | 2 +- 3 files changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp index aedadcf77db39..6ea27695955d7 100644 --- a/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp +++ b/sycl/test-e2e/Plugin/interop-level-zero-image-get-native-mem.cpp @@ -103,4 +103,4 @@ int main() { std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; #endif return 0; -} \ No newline at end of file +} diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp index 329ac64a47dbd..2d1a57df78447 100644 --- a/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp +++ b/sycl/test-e2e/Plugin/interop-level-zero-image-ownership.cpp @@ -134,4 +134,4 @@ int main() { #endif std::cout << "chau" << std::endl; return 0; -} \ No newline at end of file +} diff --git a/sycl/test-e2e/Plugin/interop-level-zero-image.cpp b/sycl/test-e2e/Plugin/interop-level-zero-image.cpp index 4c392b1e03582..fe87e730d6b67 100644 --- a/sycl/test-e2e/Plugin/interop-level-zero-image.cpp +++ b/sycl/test-e2e/Plugin/interop-level-zero-image.cpp @@ -217,4 +217,4 @@ int main() { std::cout << "Missing Level-Zero backend. Test skipped." << std::endl; #endif return 0; -} \ No newline at end of file +}