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 2b64dd1da2cf3..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 @@ -17,6 +17,7 @@ NOTE: By necessity, this specification exposes some details about the way SYCL i |2|Added support for the make_buffer() API. |3|Added device member to backend_input_t. |4|Change the definition of backend_input_t and backend_return_t for the queue object, which changes the API for make_queue and get_native (when applied to queue). +|5|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. @@ -43,15 +44,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 @@ -247,6 +248,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; + sycl::range Range; + ext::oneapi::level_zero::ownership Ownership{ + ext::oneapi::level_zero::ownership::transfer}; + } +``` + + ### 4.2 Obtaining of native Level-Zero handles from SYCL objects @@ -264,7 +287,7 @@ It is currently supported for SYCL ```platform```, ```device```, ```context```, The ```get_native(queue)``` function returns either ```ze_command_queue_handle_t``` or ```ze_command_list_handle_t``` depending on the manner in which the input argument ```queue``` had been created. Queues created with the SYCL ```queue``` constructors have a default setting for whether they use command queues or command lists. The default and how it may be changed is documented in the description for the environment variable ```SYCL_PI_LEVEL_ZERO_USE_IMMEDIATE_COMMANDLISTS```. Queues created using ```make_queue()``` use either a command list or command queue depending on the input argument to ```make_queue``` and are not affected by the default for SYCL queues or the environment variable. 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``` @@ -299,9 +322,9 @@ an application to create a SYCL object that encapsulates a corresponding Level-Z ``` C++ -make_platform( - const backend_input_t< - backend::ext_oneapi_level_zero, platform> &) +template +platform make_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. @@ -309,9 +332,9 @@ make_platform( ``` C++ -make_device( - const backend_input_t< - backend::ext_oneapi_level_zero, device> &) +template +device make_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. @@ -319,9 +342,9 @@ make_device( ``` C++ -make_context( - const backend_input_t< - backend::ext_oneapi_level_zero, context> &) +template +context make_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. @@ -329,9 +352,9 @@ make_context( ``` C++ -make_queue( - const backend_input_t< - backend::ext_oneapi_level_zero, queue> &, +template +queue make_queue( + const backend_input_t &, const context &Context) ``` @@ -346,9 +369,9 @@ the ```compute_index``` property which is built into the command queue or comman ``` C++ -make_event( - const backend_input_t< - backend::ext_oneapi_level_zero, event> &, +template +event make_event( + const backend_input_t &, const context &Context) ``` @@ -357,11 +380,11 @@ make_event( ``` C++ -make_kernel_bundle( - const backend_input_t< - backend::ext_oneapi_level_zero, - kernel_bundle> &, +// State must be bundle_state::executable +template +kernel_bundle make_kernel_bundle( + const backend_input_t> &, const context &Context) ``` @@ -383,9 +406,9 @@ interoperability kernel_bundle destructor is called. ``` C++ -make_kernel( - const backend_input_t< - backend::ext_oneapi_level_zero, kernel> &, +template +kernel make_kernel( + const backend_input_t &, const context &Context) ``` @@ -405,9 +428,15 @@ Level-Zero kernel ``` C++ -make_buffer( - const backend_input_t> &, +template >> +buffer make_buffer( + const backend_input_t> &, const context &Context) ``` @@ -421,9 +450,15 @@ Synchronization rules for a buffer that is created with this API are described i ``` C++ -make_buffer( - const backend_input_t> &, +template >> +buffer make_buffer( + const backend_input_t> &, const context &Context, event AvailableEvent) ``` @@ -433,6 +468,94 @@ 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++ +template +image make_image( + const backend_input_t> &backendObject, + const context &targetContext); +``` + +This API is available starting with revision 5 of this specification. + +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}; + }; +``` +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++ +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; +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; +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); +``` + +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. + + + + + + +``` C++ +template +image make_image( + const backend_input_t> &backendObject, + const context &targetContext, event availableEvent); +``` + +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. +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. + + NOTE: We shall consider adding other interoperability as needed, if possible. @@ -509,4 +632,5 @@ 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 -|10|2023-03-14|Rajiv Deodhar|Added support for Level Zero immediate command lists \ No newline at end of file +|11|2023-03-14|Rajiv Deodhar|Added support for Level Zero immediate command lists +|12|2023-04-06|Chris Perkins|Introduced make_image() API \ No newline at end of file diff --git a/sycl/include/sycl/backend.hpp b/sycl/include/sycl/backend.hpp index 98278071628b2..a70a9ef6ba657 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 @@ -335,6 +336,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 412144d77a996..b68df391ae56f 100644 --- a/sycl/include/sycl/detail/backend_traits_level_zero.hpp +++ b/sycl/include/sycl/detail/backend_traits_level_zero.hpp @@ -165,6 +165,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 = std::variant; @@ -214,6 +235,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 02c2f38cf3b38..fd5bc5a844cef 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(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 1a3a42cd7a6b1..38d173137c82e 100644 --- a/sycl/include/sycl/detail/pi.h +++ b/sycl/include/sycl/detail/pi.h @@ -86,9 +86,11 @@ // 12.27 Added new queue create and get APIs for immediate commandlists // piextQueueCreate2, piextQueueCreateWithNativeHandle2, // piextQueueGetNativeHandle2 +// 12.28 Added piextMemImageCreateWithNativeHandle for creating images from +// native handles. #define _PI_H_VERSION_MAJOR 12 -#define _PI_H_VERSION_MINOR 27 +#define _PI_H_VERSION_MINOR 28 #define _PI_STRING_HELPER(a) #a #define _PI_CONCAT(a, b) _PI_STRING_HELPER(a.b) @@ -1308,6 +1310,24 @@ __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. +/// +/// \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 piextMemImageCreateWithNativeHandle( + 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 cd322e1cc2ebd..9ab4c0ed63d8f 100644 --- a/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp +++ b/sycl/include/sycl/ext/oneapi/backend/level_zero.hpp @@ -229,6 +229,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..5dad481ead77a 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,31 @@ 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 + 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 +188,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 4c40d0f78c4ab..68b4fb8ef60b9 100644 --- a/sycl/plugins/cuda/pi_cuda.cpp +++ b/sycl/plugins/cuda/pi_cuda.cpp @@ -2452,6 +2452,29 @@ 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] 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_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 {}; +} + /// 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 b6796fe7e689f..65b90cc8db3b5 100644 --- a/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp +++ b/sycl/plugins/esimd_emulator/pi_esimd_emulator.cpp @@ -1314,6 +1314,12 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle, pi_context, bool, DIE_NO_IMPLEMENTATION; } +pi_result piextMemImageCreateWithNativeHandle(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 625898f04ed02..311af5da41ab4 100644 --- a/sycl/plugins/hip/pi_hip.cpp +++ b/sycl/plugins/hip/pi_hip.cpp @@ -2451,6 +2451,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[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 +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) { + (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 4967a46489f26..d37e558fbe03f 100644 --- a/sycl/plugins/level_zero/pi_level_zero.cpp +++ b/sycl/plugins/level_zero/pi_level_zero.cpp @@ -3184,8 +3184,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); @@ -3254,12 +3256,15 @@ pi_result piMemRelease(pi_mem Mem) { if (Mem->isImage()) { char *ZeHandleImage; - PI_CALL(Mem->getZeHandle(ZeHandleImage, _pi_mem::write_only)); - 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); + 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))); + // Gracefully handle the case that L0 was already unloaded. + if (ZeResult && ZeResult != ZE_RESULT_ERROR_UNINITIALIZED) + return mapError(ZeResult); + } } else { auto Buffer = static_cast(Mem); Buffer->free(); @@ -3269,20 +3274,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); - +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) { @@ -3393,8 +3387,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); @@ -3403,6 +3397,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 = + pi2zeImageDesc(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 @@ -3416,7 +3433,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 @@ -3545,6 +3562,42 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle NativeHandle, return PI_SUCCESS; } +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); + 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 = + pi2zeImageDesc(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 bf193f6b761e6..4d55fe16a0374 100644 --- a/sycl/plugins/level_zero/pi_level_zero.hpp +++ b/sycl/plugins/level_zero/pi_level_zero.hpp @@ -992,11 +992,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 { @@ -1018,6 +1021,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 414c4b3fc281b..48fd7dc5017db 100644 --- a/sycl/plugins/opencl/pi_opencl.cpp +++ b/sycl/plugins/opencl/pi_opencl.cpp @@ -1212,6 +1212,19 @@ pi_result piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } +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; + (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..95022c60d72ce 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/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index 89c4fb41507bd..452af338ea483 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -66,7 +66,7 @@ __SYCL_INLINE_VER_NAMESPACE(_V1) { #define SYCL_EXT_INTEL_USM_ADDRESS_SPACES 2 #define SYCL_EXT_INTEL_RUNTIME_BUFFER_LOCATION 1 #define SYCL_EXT_INTEL_QUEUE_INDEX 1 -#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 4 +#define SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO 5 #define SYCL_EXT_ONEAPI_USM_DEVICE_READ_ONLY 1 #define SYCL_EXT_ONEAPI_KERNEL_PROPERTIES 1 #define SYCL_EXT_ONEAPI_QUEUE_EMPTY 1 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/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..6ea27695955d7 --- /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; +} 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..2d1a57df78447 --- /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; +} 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..fe87e730d6b67 --- /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; +} diff --git a/sycl/test/abi/pi_level_zero_symbol_check.dump b/sycl/test/abi/pi_level_zero_symbol_check.dump index c47d8d9f428ae..08292e4fbeb45 100644 --- a/sycl/test/abi/pi_level_zero_symbol_check.dump +++ b/sycl/test/abi/pi_level_zero_symbol_check.dump @@ -100,6 +100,7 @@ piextKernelSetArgPointer piextKernelSetArgSampler piextMemCreateWithNativeHandle piextMemGetNativeHandle +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 f62f5e6ca59ac..a367ec1afd33d 100644 --- a/sycl/test/abi/pi_opencl_symbol_check.dump +++ b/sycl/test/abi/pi_opencl_symbol_check.dump @@ -48,6 +48,7 @@ piextKernelSetArgPointer piextKernelSetArgSampler piextMemCreateWithNativeHandle piextMemGetNativeHandle +piextMemImageCreateWithNativeHandle piextPlatformCreateWithNativeHandle piextPlatformGetNativeHandle piextProgramCreateWithNativeHandle diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index cc05b7a3ee3c6..4390b819db62f 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3734,7 +3734,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 @@ -3745,8 +3747,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 @@ -3762,6 +3766,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 @@ -3770,6 +3775,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 diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 75c22f1d7a0c1..5f56778c2c322 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -381,6 +381,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 @@ -488,6 +489,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 @@ -496,6 +498,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 diff --git a/sycl/test/extensions/macro.cpp b/sycl/test/extensions/macro.cpp index 15c90af64fca6..b6a3efed81c91 100644 --- a/sycl/test/extensions/macro.cpp +++ b/sycl/test/extensions/macro.cpp @@ -15,7 +15,7 @@ constexpr bool sub_group_mask_macro_defined = true; constexpr bool sub_group_mask_macro_defined = false; #endif -#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO == 4 +#if SYCL_EXT_ONEAPI_BACKEND_LEVEL_ZERO == 5 constexpr bool backend_level_zero_macro_defined = true; #else constexpr bool backend_level_zero_macro_defined = false; diff --git a/sycl/unittests/helpers/PiMockPlugin.hpp b/sycl/unittests/helpers/PiMockPlugin.hpp index 7ebab2627b506..a2fd198c4055d 100644 --- a/sycl/unittests/helpers/PiMockPlugin.hpp +++ b/sycl/unittests/helpers/PiMockPlugin.hpp @@ -498,6 +498,15 @@ mock_piextMemCreateWithNativeHandle(pi_native_handle nativeHandle, return PI_SUCCESS; } +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) { + *RetImage = reinterpret_cast(NativeHandle); + retainDummyHandle(*RetImage); + return PI_SUCCESS; +} + // // Program //