From b1aab046f40880ededbc7cbbef96a4817ca2cf31 Mon Sep 17 00:00:00 2001 From: przemektmalon Date: Tue, 25 Jul 2023 17:24:37 +0100 Subject: [PATCH] [SYCL][Bindless][3/4] Add experimental implementation of SYCL bindless images extension (#10454) This commit stands as the third commit of four to make code review easier, mostly covering the changes made to the user-facing SYCL API for the [bindless images extension proposal](https://github.com/intel/llvm/pull/9842). ### Overview The bindless images extension provides a new interface for allocating, creating, and accessing images in SYCL. Image memory allocation is seperated from image handle creation, and image handles can be passed to kernels without requesting access through accessors. This approach provides much more flexibility to the user, as well as enabling programs to implement features that were impossible to implement using standard SYCL images, such as a texture atlas. In addition to providing a new interface for images, this extension also provides initial experimental support for importing external memory into SYCL. ### Following Split PRs - [4/4] Add tests ### Authors Co-authored-by: Isaac Ault Co-authored-by: Hugh Bird Co-authored-by: Duncan Brawley Co-authored-by: Przemek Malon Co-authored-by: Chedy Najjar Co-authored-by: Sean Stirling Co-authored-by: Peter Zuzek --- .../llvm/SYCLLowerIR/DeviceConfigFile.td | 17 +- sycl/include/sycl/detail/cg.hpp | 74 ++ sycl/include/sycl/device_aspect_macros.hpp | 110 +++ .../sycl/ext/oneapi/bindless_images.hpp | 772 ++++++++++++++++++ .../ext/oneapi/bindless_images_descriptor.hpp | 90 ++ .../ext/oneapi/bindless_images_interop.hpp | 56 ++ .../ext/oneapi/bindless_images_memory.hpp | 119 +++ .../ext/oneapi/bindless_images_sampler.hpp | 47 ++ sycl/include/sycl/handler.hpp | 144 ++++ sycl/include/sycl/info/aspects.def | 11 + .../sycl/info/ext_oneapi_device_traits.def | 19 + sycl/include/sycl/queue.hpp | 677 +++++++++++++++ sycl/include/sycl/sycl.hpp | 1 + sycl/source/CMakeLists.txt | 1 + sycl/source/detail/bindless_images.cpp | 750 +++++++++++++++++ sycl/source/detail/device_impl.cpp | 89 ++ sycl/source/detail/device_info.hpp | 39 + sycl/source/detail/graph_impl.hpp | 6 + sycl/source/detail/handler_impl.hpp | 13 + sycl/source/detail/memory_manager.cpp | 31 + sycl/source/detail/memory_manager.hpp | 12 + sycl/source/detail/scheduler/commands.cpp | 40 + sycl/source/feature_test.hpp.in | 1 + sycl/source/handler.cpp | 257 ++++++ sycl/test/abi/sycl_symbols_linux.dump | 77 ++ sycl/test/abi/sycl_symbols_windows.dump | 109 +++ 26 files changed, 3561 insertions(+), 1 deletion(-) create mode 100644 sycl/include/sycl/ext/oneapi/bindless_images.hpp create mode 100644 sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp create mode 100644 sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp create mode 100644 sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp create mode 100644 sycl/include/sycl/ext/oneapi/bindless_images_sampler.hpp create mode 100644 sycl/source/detail/bindless_images.cpp diff --git a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td index 2ca3ed85676cb..c4690e418e9b9 100644 --- a/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td +++ b/llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td @@ -50,6 +50,17 @@ def AspectExt_intel_memory_clock_rate : Aspect<"ext_intel_memory_clock_rate">; def AspectExt_intel_memory_bus_width : Aspect<"ext_intel_memory_bus_width">; def AspectEmulated : Aspect<"emulated">; def AspectExt_intel_legacy_image : Aspect<"ext_intel_legacy_image">; +def AspectExt_oneapi_bindless_images : Aspect<"ext_oneapi_bindless_images">; +def AspectExt_oneapi_bindless_images_shared_usm : Aspect<"ext_oneapi_bindless_images_shared_usm">; +def AspectExt_oneapi_bindless_images_1d_usm : Aspect<"ext_oneapi_bindless_images_1d_usm">; +def AspectExt_oneapi_bindless_images_2d_usm : Aspect<"ext_oneapi_bindless_images_2d_usm">; +def AspectExt_oneapi_interop_memory_import : Aspect<"ext_oneapi_interop_memory_import">; +def AspectExt_oneapi_interop_memory_export : Aspect<"ext_oneapi_interop_memory_export">; +def AspectExt_oneapi_interop_semaphore_import : Aspect<"ext_oneapi_interop_semaphore_import">; +def AspectExt_oneapi_interop_semaphore_export : Aspect<"ext_oneapi_interop_semaphore_export">; +def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">; +def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">; +def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">; // Deprecated aspects def AspectInt64_base_atomics : Aspect<"int64_base_atomics">; def AspectInt64_extended_atomics : Aspect<"int64_extended_atomics">; @@ -94,7 +105,11 @@ def : TargetInfo<"__TestAspectList", AspectExt_oneapi_native_assert, AspectHost_debuggable, AspectExt_intel_gpu_hw_threads_per_eu, AspectExt_oneapi_cuda_async_barrier, AspectExt_oneapi_bfloat16_math_functions, AspectExt_intel_free_memory, AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectEmulated, - AspectExt_intel_legacy_image], + AspectExt_intel_legacy_image, AspectExt_oneapi_bindless_images, + AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm, + AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export, + AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export, + AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference], []>; // This definition serves the only purpose of testing whether the deprecated aspect list defined in here and in SYCL RT // match. diff --git a/sycl/include/sycl/detail/cg.hpp b/sycl/include/sycl/detail/cg.hpp index b29be33cb0be7..68197e4ce599a 100644 --- a/sycl/include/sycl/detail/cg.hpp +++ b/sycl/include/sycl/detail/cg.hpp @@ -73,6 +73,9 @@ class CG { CopyFromDeviceGlobal = 20, ReadWriteHostPipe = 21, ExecCommandBuffer = 22, + CopyImage = 23, + SemaphoreWait = 24, + SemaphoreSignal = 25, }; struct StorageInitHelper { @@ -496,6 +499,77 @@ class CGCopyFromDeviceGlobal : public CG { size_t getNumBytes() { return MNumBytes; } size_t getOffset() { return MOffset; } }; +/// "Copy Image" command group class. +class CGCopyImage : public CG { + void *MSrc; + void *MDst; + sycl::detail::pi::PiMemImageDesc MImageDesc; + sycl::detail::pi::PiMemImageFormat MImageFormat; + sycl::detail::pi::PiImageCopyFlags MImageCopyFlags; + sycl::detail::pi::PiImageOffset MSrcOffset; + sycl::detail::pi::PiImageOffset MDstOffset; + sycl::detail::pi::PiImageRegion MHostExtent; + sycl::detail::pi::PiImageRegion MCopyExtent; + +public: + CGCopyImage(void *Src, void *Dst, sycl::detail::pi::PiMemImageDesc ImageDesc, + sycl::detail::pi::PiMemImageFormat ImageFormat, + sycl::detail::pi::PiImageCopyFlags ImageCopyFlags, + sycl::detail::pi::PiImageOffset SrcOffset, + sycl::detail::pi::PiImageOffset DstOffset, + sycl::detail::pi::PiImageRegion HostExtent, + sycl::detail::pi::PiImageRegion CopyExtent, + CG::StorageInitHelper CGData, detail::code_location loc = {}) + : CG(CopyImage, std::move(CGData), std::move(loc)), MSrc(Src), MDst(Dst), + MImageDesc(ImageDesc), MImageFormat(ImageFormat), + MImageCopyFlags(ImageCopyFlags), MSrcOffset(SrcOffset), + MDstOffset(DstOffset), MHostExtent(HostExtent), + MCopyExtent(CopyExtent) {} + + void *getSrc() const { return MSrc; } + void *getDst() const { return MDst; } + sycl::detail::pi::PiMemImageDesc getDesc() const { return MImageDesc; } + sycl::detail::pi::PiMemImageFormat getFormat() const { return MImageFormat; } + sycl::detail::pi::PiImageCopyFlags getCopyFlags() const { + return MImageCopyFlags; + } + sycl::detail::pi::PiImageOffset getSrcOffset() const { return MSrcOffset; } + sycl::detail::pi::PiImageOffset getDstOffset() const { return MDstOffset; } + sycl::detail::pi::PiImageRegion getHostExtent() const { return MHostExtent; } + sycl::detail::pi::PiImageRegion getCopyExtent() const { return MCopyExtent; } +}; + +/// "Semaphore Wait" command group class. +class CGSemaphoreWait : public CG { + sycl::detail::pi::PiInteropSemaphoreHandle MInteropSemaphoreHandle; + +public: + CGSemaphoreWait( + sycl::detail::pi::PiInteropSemaphoreHandle InteropSemaphoreHandle, + CG::StorageInitHelper CGData, detail::code_location loc = {}) + : CG(SemaphoreWait, std::move(CGData), std::move(loc)), + MInteropSemaphoreHandle(InteropSemaphoreHandle) {} + + sycl::detail::pi::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const { + return MInteropSemaphoreHandle; + } +}; + +/// "Semaphore Signal" command group class. +class CGSemaphoreSignal : public CG { + sycl::detail::pi::PiInteropSemaphoreHandle MInteropSemaphoreHandle; + +public: + CGSemaphoreSignal( + sycl::detail::pi::PiInteropSemaphoreHandle InteropSemaphoreHandle, + CG::StorageInitHelper CGData, detail::code_location loc = {}) + : CG(SemaphoreSignal, std::move(CGData), std::move(loc)), + MInteropSemaphoreHandle(InteropSemaphoreHandle) {} + + sycl::detail::pi::PiInteropSemaphoreHandle getInteropSemaphoreHandle() const { + return MInteropSemaphoreHandle; + } +}; /// "Execute command-buffer" command group class. class CGExecCommandBuffer : public CG { diff --git a/sycl/include/sycl/device_aspect_macros.hpp b/sycl/include/sycl/device_aspect_macros.hpp index 4c2c5dcbada4a..d42bad63427af 100644 --- a/sycl/include/sycl/device_aspect_macros.hpp +++ b/sycl/include/sycl/device_aspect_macros.hpp @@ -218,6 +218,61 @@ #define __SYCL_ALL_DEVICES_HAVE_ext_intel_legacy_image__ 0 #endif +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images__ +// __SYCL_ASPECT(ext_oneapi_bindless_images, 42) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_shared_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_shared_usm, 43) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_shared_usm__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_1d_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_1d_usm, 44) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_1d_usm__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_2d_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_2d_usm, 45) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_2d_usm__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_import__ +//__SYCL_ASPECT(ext_oneapi_interop_memory_import, 46) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_import__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__ +//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__ +//__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__ +//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__ +//__SYCL_ASPECT(ext_oneapi_mipmap, 50) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_anisotropy__ +//__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_anisotropy__ 0 +#endif + +#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_level_reference__ +//__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52) +#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap_level_reference__ 0 +#endif + #ifndef __SYCL_ANY_DEVICE_HAS_host__ // __SYCL_ASPECT(host, 0) #define __SYCL_ANY_DEVICE_HAS_host__ 0 @@ -427,3 +482,58 @@ // __SYCL_ASPECT(ext_intel_legacy_image, 41) #define __SYCL_ANY_DEVICE_HAS_ext_intel_legacy_image__ 0 #endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images__ +// __SYCL_ASPECT(ext_oneapi_bindless_images, 42) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_shared_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_shared_usm, 43) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_shared_usm__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_1d_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_1d_usm, 44) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_1d_usm__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_2d_usm__ +//__SYCL_ASPECT(ext_oneapi_bindless_images_2d_usm, 45) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_2d_usm__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_import__ +//__SYCL_ASPECT(ext_oneapi_interop_memory_import, 46) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_import__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__ +//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__ +//__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__ +//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__ +//__SYCL_ASPECT(ext_oneapi_mipmap, 50) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_anisotropy__ +//__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_anisotropy__ 0 +#endif + +#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_level_reference__ +//__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52) +#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap_level_reference__ 0 +#endif diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp new file mode 100644 index 0000000000000..4b56fb3d7583b --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -0,0 +1,772 @@ +//==----------- bindless_images.hpp --- SYCL bindless images ---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +/// Opaque unsampled image handle type. +struct unsampled_image_handle { + using raw_handle_type = pi_uint64; + raw_handle_type raw_handle; +}; +/// Opaque sampled image handle type. +struct sampled_image_handle { + using raw_handle_type = pi_uint64; + raw_handle_type raw_handle; +}; + +/** + * @brief Allocate image memory based on image_descriptor + * + * @param desc The image descriptor + * @param syclDevice The device in which we create our memory handle + * @param syclContext The context in which we create our memory handle + * @return Memory handle to allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle +alloc_image_mem(const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Allocate image memory based on image_descriptor + * + * @param desc The image descriptor + * @param syclQueue The queue in which we create our memory handle + * @return Memory handle to allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle alloc_image_mem(const image_descriptor &desc, + const sycl::queue &syclQueue); + +/** + * @brief Free image memory + * + * @param handle Memory handle to allocated memory on the device + * @param syclDevice The device in which we create our memory handle + * @param syclContext The context in which we created our memory handle + */ +__SYCL_EXPORT void free_image_mem(image_mem_handle handle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Free image memory + * + * @param handle Memory handle to allocated memory on the device + * @param syclQueue The queue in which we create our memory handle + */ +__SYCL_EXPORT void free_image_mem(image_mem_handle handle, + const sycl::device &syclQueue); + +/** + * @brief Allocate mipmap memory based on image_descriptor + * + * @param desc The image descriptor + * @param syclDevice The device in which we create our memory handle + * @param syclContext The context in which we create our memory handle + * @return Memory handle to allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle +alloc_mipmap_mem(const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Allocate mipmap memory based on image_descriptor + * + * @param desc The image descriptor + * @param syclQueue The queue in which we create our memory handle + * @return Memory handle to allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle alloc_mipmap_mem(const image_descriptor &desc, + const sycl::device &syclQueue); + +/** + * @brief Free mipmap memory + * + * @param handle The mipmap memory handle + * @param syclDevice The device in which we created our memory handle + * @param syclContext The context in which we created our memory handle + */ +__SYCL_EXPORT void free_mipmap_mem(image_mem_handle handle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Free mipmap memory + * + * @param handle The mipmap memory handle + * @param syclQueue The queue in which we created our memory handle + */ +__SYCL_EXPORT void free_mipmap_mem(image_mem_handle handle, + const sycl::queue &syclQueue); + +/** + * @brief Retrieve the memory handle to an individual mipmap image + * + * @param mipMem The memory handle to the mipmapped array + * @param level The requested level of the mipmap + * @param syclDevice The device in which we created our memory handle + * @param syclContext The context in which we created our memory handle + * @return Memory handle to the individual mipmap image + */ +__SYCL_EXPORT image_mem_handle get_mip_level_mem_handle( + const image_mem_handle mipMem, const unsigned int level, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Retrieve the memory handle to an individual mipmap image + * + * @param mipMem The memory handle to the mipmapped array + * @param level The requested level of the mipmap + * @param syclQueue The queue in which we created our memory handle + * @return Memory handle to the individual mipmap image + */ +__SYCL_EXPORT image_mem_handle get_mip_level_mem_handle( + const image_mem_handle mipMem, const unsigned int level, + const sycl::device &syclQueue); + +/** + * @brief Import external memory taking an external memory handle (the type + * of which is dependent on the OS & external API) and return an + * interop memory handle + * + * @tparam ExternalMemHandleType Handle type describing external memory handle + * @param externalMem External memory descriptor + * @param syclDevice The device in which we create our interop memory + * @param syclContext The context in which we create our interop memory + * handle + * @return Interop memory handle to the external memory + */ +template +__SYCL_EXPORT interop_mem_handle import_external_memory( + external_mem_descriptor externalMem, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Import external memory taking an external memory handle (the type + * of which is dependent on the OS & external API) and return an + * interop memory handle + * + * @tparam ExternalMemHandleType Handle type describing external memory handle + * @param externalMem External memory descriptor + * @param syclQueue The queue in which we create our interop memory + * handle + * @return Interop memory handle to the external memory + */ +template +__SYCL_EXPORT interop_mem_handle import_external_memory( + external_mem_descriptor externalMem, + const sycl::queue &syclQueue); + +/** + * @brief Maps an interop memory handle to an image memory handle (which may + * have a device optimized memory layout) + * + * @param memHandle Interop memory handle + * @param desc The image descriptor + * @param syclDevice The device in which we create our image memory handle + * @param syclContext The conext in which we create our image memory handle + * @return Memory handle to externally allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle map_external_memory_array( + interop_mem_handle memHandle, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Maps an interop memory handle to an image memory handle (which may + * have a device optimized memory layout) + * + * @param memHandle Interop memory handle + * @param desc The image descriptor + * @param syclQueue The queue in which we create our image memory handle + * @return Memory handle to externally allocated memory on the device + */ +__SYCL_EXPORT image_mem_handle map_external_memory_array( + interop_mem_handle memHandle, const image_descriptor &descm, + const sycl::queue &syclQueue); + +/** + * @brief Import external semaphore taking an external semaphore handle (the + * type of which is dependent on the OS & external API) + * + * @tparam ExternalSemaphoreHandleType Handle type describing external + * semaphore handle + * @param externalSemaphoreDesc External semaphore descriptor + * @param syclDevice The device in which we create our interop semaphore + * handle + * @param syclContext The context in which we create our interop semaphore + * handle + * @return Interop semaphore handle to the external semaphore + */ +template +__SYCL_EXPORT interop_semaphore_handle import_external_semaphore( + external_semaphore_descriptor + externalSemaphoreDesc, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Import external semaphore taking an external semaphore handle (the + * type of which is dependent on the OS & external API) + * + * @tparam ExternalSemaphoreHandleType Handle type describing external + * semaphore handle + * @param externalSemaphoreDesc External semaphore descriptor + * @param syclQueue The queue in which we create our interop semaphore + * handle + * @return Interop semaphore handle to the external semaphore + */ +template +__SYCL_EXPORT interop_semaphore_handle import_external_semaphore( + external_semaphore_descriptor + externalSemaphoreDesc, + const sycl::queue &syclQueue); + +/** + * @brief Destroy the external semaphore handle + * + * @param semaphoreHandle The interop semaphore handle to destroy + * @param syclDevice The device in which the interop semaphore handle was + * created + * @param syclContext The context in which the interop semaphore handle was + * created + */ +__SYCL_EXPORT void +destroy_external_semaphore(interop_semaphore_handle semaphoreHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Destroy the external semaphore handle + * + * @param semaphoreHandle The interop semaphore handle to destroy + * @param syclQueue The queue in which the interop semaphore handle was + * created + */ +__SYCL_EXPORT void +destroy_external_semaphore(interop_semaphore_handle semaphoreHandle, + const sycl::queue &syclQueue); + +/** + * @brief Release external memory + * + * @param interopHandle The interop memory handle to release + * @param syclDevice The device in which the interop memory handle was + * created + * @param syclContext The context in which the interop memory handle was + * created + */ +__SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Release external memory + * + * @param interopHandle The interop memory handle to release + * @param syclQueue The queue in which the interop memory handle was + * created + */ +__SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle, + const sycl::queue &syclQueue); + +/** + * @brief Create an image and return the device image handle + * + * @param memHandle Device memory handle wrapper for allocated image memory + * @param desc The image descriptor + * @param syclDevice The device in which we created our image handle + * @param syclContext The context in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem &memHandle, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Create an image and return the device image handle + * + * @param memHandle Device memory handle wrapper for allocated image memory + * @param desc The image descriptor + * @param syclqueue The queue in which we created our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT unsampled_image_handle create_image(image_mem &memHandle, + const image_descriptor &desc, + const sycl::queue &syclQueue); + +/** + * @brief Create an image and return the device image handle + * + * @param memHandle Device memory handle for allocated image memory + * @param desc The image descriptor + * @param syclDevice The device in which we created our image handle + * @param syclContext The context in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem_handle memHandle, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext); + +/** + * @brief Create an image and return the device image handle + * + * @param memHandle Device memory handle for allocated image memory + * @param desc The image descriptor + * @param syclQueue The queue in which we created our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT unsampled_image_handle create_image(image_mem_handle memHandle, + const image_descriptor &desc, + const sycl::queue &syclQueue); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param imgMem Device memory pointer to allocated image memory + * @param pitch The allocation pitch value + * @param sampler bindless image sampler to sample the image + * @param desc The image descriptor + * @param syclDevice The device in which we create our image handle + * @param syclContext The context in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param imgMem Device memory pointer to allocated image memory + * @param pitch The allocation pitch value + * @param sampler bindless image sampler used to sample the image + * @param desc The image descriptor + * @param syclQueue The queue in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param memHandle Device memory handle wrapper for allocated image memory + * @param sampler bindless image sampler used to sample the image + * @param desc The image descriptor + * @param syclDevice The device in which we create our image handle + * @param syclContext The context in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(image_mem &memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param memHandle Device memory handle wrapper for allocated image memory + * @param sampler bindless image sampler used to sample the image + * @param desc The image descriptor + * @param syclQueue The queue in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(image_mem &memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param memHandle Device memory handle for allocated image memory + * @param sampler bindless image sampler used to sample the image + * @param desc The image descriptor + * @param syclDevice The device in which we create our image handle + * @param syclContext The context in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Create a sampled image and return the device image handle + * + * @param memHandle Device memory handle for allocated image memory + * @param sampler bindless image sampler used to sample the image + * @param desc The image descriptor + * @param syclQueue The queue in which we create our image handle + * @return Image handle to created image object on the device + */ +__SYCL_EXPORT sampled_image_handle +create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue); + +/** + * @brief Destroy an unsampled image handle. Does not free memory backing the + * handle + * + * @param imageHandle The unsampled image handle to destroy + * @param syclDevice The device in which we created our image handle + * @param syclContext The context in which we created our image handle + **/ +__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Destroy an unsampled image handle. Does not free memory backing the + * handle + * + * @param imageHandle The unsampled image handle to destroy + * @param syclQueue The queue in which we created our image handle + **/ +__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle, + const sycl::queue &syclQueue); + +/** + * @brief Destroy a sampled image handle. Does not free memory backing the + * handle + * + * @param imageHandle The sampled image handle to destroy + * @param syclDevice The device in which we created our image handle + * @param syclContext The context in which we created our image handle + **/ +__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Destroy a sampled image handle. Does not free memory backing the + * handle + * + * @param imageHandle The sampled image handle to destroy + * @param syclQueue The queue in which we created our image handle + **/ +__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle, + const sycl::queue &syclQueue); + +/** + * @brief Allocate pitched USM image memory + * + * @param resultPitch The allocation pitch value + * @param widthInBytes The width of the image in bytes + * @param height The height of the image in elements + * @param elementSizeBytes Number of bytes of a singular image element + * @param syclQueue The queue + * @return Generic pointer to allocated USM image memory + */ +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + size_t widthInBytes, size_t height, + unsigned int elementSizeBytes, + const sycl::queue &syclQueue); + +/** + * @brief Allocate pitched USM image memory + * + * @param resultPitch The allocation pitch value + * @param widthInBytes The width of the image in bytes + * @param height The height of the image in elements + * @param elementSizeBytes Number of bytes of a singular image element + * @param syclDevice The device + * @param syclContext The context + * @return Generic pointer to allocated USM image memory + */ +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + size_t widthInBytes, size_t height, + unsigned int elementSizeBytes, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Allocate pitched USM image memory + * + * @param resultPitch The allocation pitch value + * @param desc The image descriptor + * @param syclQueue The queue + * @return Generic pointer to allocated USM image memory + */ +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + const image_descriptor &desc, + const sycl::queue &syclQueue); + +/** + * @brief Allocate pitched USM image memory + * + * @param resultPitch The allocation pitch value + * @param desc The image descriptor + * @param syclDevice The device + * @param syclContext The context + * @return Generic pointer to allocated USM image memory + */ +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + const image_descriptor &desc, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Get the range that describes the image's dimensions + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclDevice The device in which we created our image memory handle + * @param syclContext The context in which we created our image memory handle + * @return sycl range describing image's dimensions + */ +__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Get the range that describes the image's dimensions + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclQueue The queue in which we created our image memory handle + * @return sycl range describing image's dimensions + */ +__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle, + const sycl::queue &syclQueue); + +/** + * @brief Get the channel type that describes the image memory + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclDevice The device in which we created our image memory handle + * @param syclContext The context in which we created our image memory handle + * @return sycl image channel type that describes the image + */ +__SYCL_EXPORT sycl::image_channel_type +get_image_channel_type(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Get the channel type that describes the image memory + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclQueue The queue in which we created our image memory handle + * @return sycl image channel type that describes the image + */ +__SYCL_EXPORT sycl::image_channel_type +get_image_channel_type(const image_mem_handle memHandle, + const sycl::queue &syclQueue); + +/** + * @brief Get the number of channels that describes the image memory + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclDevice The device in which we created our image memory handle + * @param syclContext The context in which we created our image memory handle + * @return The number of channels describing the image + */ +__SYCL_EXPORT unsigned int +get_image_num_channels(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Get the number of channels that describes the image memory + * + * @param memHandle Memory handle to allocated memory on the device + * @param syclQueue The queue in which we created our image memory handle + * @return The number of channels describing the image + */ +__SYCL_EXPORT unsigned int +get_image_num_channels(const image_mem_handle memHandle, + const sycl::queue &syclQueue); + +namespace detail { +// Get the number of coordinates +template constexpr size_t coord_size() { + if constexpr (std::is_scalar::value) { + return 1; + } else { + return CoordT::size(); + } +} +} // namespace detail + +/** + * @brief Read an unsampled image using its handle + * + * @tparam DataT The return type + * @tparam CoordT The input coordinate type. e.g. int, int2, or int4 for + * 1D, 2D, and 3D respectively + * @param imageHandle The image handle + * @param coords The coordinates at which to fetch image data + * @return Image data + * + * __NVPTX__: Name mangling info + * Cuda surfaces require integer coords (by bytes) + * Cuda textures require float coords (by element or normalized) + * The name mangling should therefore not interfere with one + * another + */ +template +DataT read_image(const unsampled_image_handle &imageHandle, + const CoordT &coords) { + constexpr size_t coordSize = detail::coord_size(); + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, + "Expected input coordinate to be have 1, 2, or 4 components " + "for 1D, 2D and 3D images respectively."); + +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__NVPTX__) + return __invoke__ImageRead(imageHandle.raw_handle, + coords); +#else + // TODO: add SPIRV part for unsampled image read +#endif +#else + assert(false); // Bindless images not yet implemented on host +#endif +} + +/** + * @brief Read a sampled image using its handle + * + * @tparam DataT The return type + * @tparam CoordT The input coordinate type. e.g. float, float2, or float4 for + * 1D, 2D, and 3D respectively + * @param imageHandle The image handle + * @param coords The coordinates at which to fetch image data + * @return Sampled image data + * + * __NVPTX__: Name mangling info + * Cuda surfaces require integer coords (by bytes) + * Cuda textures require float coords (by element or normalized) + * The name mangling should therefore not interfere with one + * another + */ +template +DataT read_image(const sampled_image_handle &imageHandle, + const CoordT &coords) { + constexpr size_t coordSize = detail::coord_size(); + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, + "Expected input coordinate to be have 1, 2, or 4 components " + "for 1D, 2D and 3D images respectively."); + +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__NVPTX__) + return __invoke__ImageRead(imageHandle.raw_handle, + coords); +#else + // TODO: add SPIRV part for sampled image read +#endif +#else + assert(false); // Bindless images not yet implemented on host. +#endif +} + +/** + * @brief Read a mipmap image using its handle with LOD filtering + * + * @tparam DataT The return type + * @tparam CoordT The input coordinate type. e.g. float, float2, or float4 for + * 1D, 2D, and 3D respectively + * @param imageHandle The mipmap image handle + * @param coords The coordinates at which to fetch mipmap image data + * @param level The mipmap level at which to sample + * @return Mipmap image data with LOD filtering + */ +template +DataT read_image(const sampled_image_handle &imageHandle, const CoordT &coords, + const float level) { + constexpr size_t coordSize = detail::coord_size(); + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, + "Expected input coordinate to be have 1, 2, or 4 components " + "for 1D, 2D and 3D images respectively."); + +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__NVPTX__) + return __invoke__ImageReadLod(imageHandle.raw_handle, + coords, level); +#else + // TODO: add SPIRV for mipmap level read +#endif +#else + assert(false); // Bindless images not yet implemented on host +#endif +} + +/** + * @brief Read a mipmap image using its handle with anisotropic filtering + * + * @tparam DataT The return type + * @tparam CoordT The input coordinate type. e.g. float, float2, or float4 for + * 1D, 2D, and 3D respectively + * @param imageHandle The mipmap image handle + * @param coords The coordinates at which to fetch mipmap image data + * @param dX Screen space gradient in the x dimension + * @param dY Screen space gradient in the y dimension + * @return Mipmap image data with anisotropic filtering + */ +template +DataT read_image(const sampled_image_handle &imageHandle, const CoordT &coords, + const CoordT &dX, const CoordT &dY) { + constexpr size_t coordSize = detail::coord_size(); + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, + "Expected input coordinate and gradient to have 1, 2, or 4 " + "components " + "for 1D, 2D and 3D images respectively."); + +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__NVPTX__) + return __invoke__ImageReadGrad( + imageHandle.raw_handle, coords, dX, dY); +#else + // TODO: add SPIRV part for mipmap grad read +#endif +#else + assert(false); // Bindless images not yet implemented on host +#endif +} + +/** + * @brief Write to an unsampled image using its handle + * + * @tparam DataT The data type to write + * @tparam CoordT The input coordinate type. e.g. int, int2, or int4 for + * 1D, 2D, and 3D respectively + * @param imageHandle The image handle + * @param coords The coordinates at which to write image data + */ +template +void write_image(const unsampled_image_handle &imageHandle, + const CoordT &Coords, const DataT &Color) { + constexpr size_t coordSize = detail::coord_size(); + static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4, + "Expected input coordinate to be have 1, 2, or 4 components " + "for 1D, 2D and 3D images respectively."); + +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__NVPTX__) + __invoke__ImageWrite( + (uint64_t)imageHandle.raw_handle, Coords, Color); +#else + // TODO: add SPIRV part for unsampled image write +#endif +#else + assert(false); // Bindless images not yet implemented on host +#endif +} + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp new file mode 100644 index 0000000000000..5576c41beeb82 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp @@ -0,0 +1,90 @@ +//==------ bindless_images_descriptor.hpp --- SYCL bindless images ---------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +/// image type enum +enum class image_type : unsigned int { + standard = 0, + interop = 1, + mipmap = 2, + cubemap = 3, /* Not implemented */ + layered = 4, /* Not implemented */ +}; + +/// A struct to describe the properties of an image. +struct image_descriptor { + size_t width; + size_t height; + size_t depth; + image_channel_order channel_order; + image_channel_type channel_type; + image_type type; + unsigned int num_levels; + + image_descriptor() = default; + + image_descriptor(range<1> dims, image_channel_order channel_order, + image_channel_type channel_type, + image_type type = image_type::standard, + unsigned int num_levels = 1) + : width(dims[0]), height(0), depth(0), channel_order(channel_order), + channel_type(channel_type), type(type), num_levels(num_levels) {} + + image_descriptor(range<2> dims, image_channel_order channel_order, + image_channel_type channel_type, + image_type type = image_type::standard, + unsigned int num_levels = 1) + : width(dims[0]), height(dims[1]), depth(0), channel_order(channel_order), + channel_type(channel_type), type(type), num_levels(num_levels) {} + + image_descriptor(range<3> dims, image_channel_order channel_order, + image_channel_type channel_type, + image_type type = image_type::standard, + unsigned int num_levels = 1) + : width(dims[0]), height(dims[1]), depth(dims[2]), + channel_order(channel_order), channel_type(channel_type), type(type), + num_levels(num_levels){}; + + /// Get the descriptor for a mipmap level + image_descriptor get_mip_level_desc(unsigned int level) const { + // Check that this descriptor describes a mipmap - otherwise throw + if (this->type != image_type::mipmap) + throw sycl::exception( + sycl::errc::invalid, + "Invalid descriptor `image_type` passed to " + "`get_mip_level_desc`. A mipmap level descriptor can only be " + "requested by a descriptor with mipmap image type!"); + + // Generate a new descriptor which represents the level accordingly + // Do not allow height/depth values to be clamped to 1 when naturally 0 + size_t width = std::max(this->width >> level, 1); + size_t height = this->height == 0 + ? this->height + : std::max(this->height >> level, 1); + size_t depth = this->depth == 0 ? this->depth + : std::max(this->depth >> level, 1); + + // This will generate the new descriptor with image_type standard + // since individual mip levels are standard images + sycl::ext::oneapi::experimental::image_descriptor levelDesc( + {width, height, depth}, this->channel_order, this->channel_type); + + return levelDesc; + } +}; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp new file mode 100644 index 0000000000000..3982a55554230 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/bindless_images_interop.hpp @@ -0,0 +1,56 @@ +//==----------- bindless_images_interop.hpp --- SYCL bindless images -------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +/// Opaque interop memory handle type +struct interop_mem_handle { + using raw_handle_type = pi_uint64; + raw_handle_type raw_handle; +}; + +/// External memory file descriptor type +struct external_mem_fd { + int file_descriptor; +}; + +/// Windows external memory type +struct external_mem_win32 { + void *handle; + const void *name; +}; + +/// Opaque external memory descriptor type +template struct external_mem_descriptor { + HandleType external_handle; + size_t size_in_bytes; +}; + +/// Opaque interop semaphore handle type +struct interop_semaphore_handle { + using raw_handle_type = pi_uint64; + raw_handle_type raw_handle; +}; + +/// External semaphore file descriptor type +struct external_semaphore_fd { + int file_descriptor; +}; + +/// Opaque external semaphore descriptor type +template struct external_semaphore_descriptor { + HandleType external_handle; +}; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp new file mode 100644 index 0000000000000..941b0315342a1 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp @@ -0,0 +1,119 @@ +//==----------- bindless_images_memory.hpp --- SYCL bindless images --------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include +#include +#include +#include + +namespace sycl { +inline namespace _V1 { + +// Forward declaration +class queue; + +namespace ext::oneapi::experimental { + +/// Opaque image memory handle type +struct image_mem_handle { + using handle_type = void *; + handle_type raw_handle; +}; + +namespace detail { + +class image_mem_impl { + using raw_handle_type = image_mem_handle; + +public: + __SYCL_EXPORT image_mem_impl(const image_descriptor &desc, + const device &syclDevice, + const context &syclContext); + __SYCL_EXPORT ~image_mem_impl(); + + raw_handle_type get_handle() const { return handle; } + const image_descriptor &get_descriptor() const { return descriptor; } + sycl::device get_device() const { return syclDevice; } + sycl::context get_context() const { return syclContext; } + +private: + raw_handle_type handle{nullptr}; + image_descriptor descriptor; + sycl::device syclDevice; + sycl::context syclContext; +}; + +} // namespace detail + +/// A class that represents image memory +class __SYCL_EXPORT image_mem { + using raw_handle_type = image_mem_handle; + +public: + image_mem() = default; + image_mem(const image_mem &) = default; + image_mem(image_mem &&rhs) = default; + + image_mem(const image_descriptor &desc, const device &syclDevice, + const context &syclContext); + image_mem(const image_descriptor &desc, const queue &syclQueue); + ~image_mem() = default; + + image_mem &operator=(const image_mem &) = default; + image_mem &operator=(image_mem &&) = default; + + bool operator==(const image_mem &rhs) const { return impl == rhs.impl; } + bool operator!=(const image_mem &rhs) const { return !(*this == rhs); } + + raw_handle_type get_handle() const { return impl->get_handle(); } + const image_descriptor &get_descriptor() const { + return impl->get_descriptor(); + } + sycl::device get_device() const { return impl->get_device(); } + sycl::context get_context() const { return impl->get_context(); } + + sycl::range<3> get_range() const; + sycl::image_channel_type get_channel_type() const; + sycl::image_channel_order get_channel_order() const; + unsigned int get_num_channels() const; + image_type get_type() const; + + raw_handle_type get_mip_level_mem_handle(const unsigned int level) const; + +protected: + std::shared_ptr impl; + + template + friend decltype(Obj::impl) + sycl::detail::getSyclObjImpl(const Obj &SyclObject); +}; + +/// Direction to copy data from bindless image handle +/// (Host -> Device) (Device -> Host) etc. +enum image_copy_flags : unsigned int { + HtoD = 0, + DtoH = 1, + DtoD = 2, +}; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl + +namespace std { +template <> struct hash { + size_t operator()( + const sycl::ext::oneapi::experimental::image_mem &image_mem) const { + return hash>()( + sycl::detail::getSyclObjImpl(image_mem)); + } +}; +} // namespace std diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_sampler.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_sampler.hpp new file mode 100644 index 0000000000000..c493f9fe199d2 --- /dev/null +++ b/sycl/include/sycl/ext/oneapi/bindless_images_sampler.hpp @@ -0,0 +1,47 @@ +//==------ bindless_images_sampler.hpp --- SYCL bindless images ------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#pragma once + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +struct bindless_image_sampler { + + bindless_image_sampler(sycl::addressing_mode addressing, + sycl::coordinate_normalization_mode coordinate, + sycl::filtering_mode filtering) + : addressing(addressing), coordinate(coordinate), filtering(filtering) {} + + bindless_image_sampler(sycl::addressing_mode addressing, + sycl::coordinate_normalization_mode coordinate, + sycl::filtering_mode filtering, + sycl::filtering_mode mipmapFiltering, + float minMipmapLevelClamp, float maxMipmapLevelClamp, + float maxAnisotropy) + : addressing(addressing), coordinate(coordinate), filtering(filtering), + mipmap_filtering(mipmapFiltering), + min_mipmap_level_clamp(minMipmapLevelClamp), + max_mipmap_level_clamp(maxMipmapLevelClamp), + max_anisotropy(maxAnisotropy) {} + + sycl::addressing_mode addressing; + sycl::coordinate_normalization_mode coordinate; + sycl::filtering_mode filtering; + sycl::filtering_mode mipmap_filtering; + float min_mipmap_level_clamp = 0.f; + float max_mipmap_level_clamp = 0.f; + float max_anisotropy = 0.f; +}; + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/include/sycl/handler.hpp b/sycl/include/sycl/handler.hpp index 7c91cbddee7ec..544a1baaffe51 100644 --- a/sycl/include/sycl/handler.hpp +++ b/sycl/include/sycl/handler.hpp @@ -41,6 +41,9 @@ #include +#include +#include + #include #include #include @@ -2889,6 +2892,147 @@ class __SYCL_EXPORT handler { ext::oneapi::experimental::graph_state::executable> Graph); + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle. An exception is + /// thrown if either \p Src is nullptr or \p Dest is incomplete. The behavior + /// is undefined if \p Desc is inconsistent with the allocated memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is an opaque image memory handle to the destination memory. + /// \param DestImgDesc is the image descriptor + void ext_oneapi_copy( + void *Src, ext::oneapi::experimental::image_mem_handle Dest, + const ext::oneapi::experimental::image_descriptor &DestImgDesc); + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. Pixel size is determined + /// by \p DestImgDesc + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an offset from the origin where the x, y, and z + /// components are measured in bytes, rows, and slices + /// respectively + /// \param SrcExtent is the extent of the source memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param Dest is an opaque image memory handle to the destination memory. + /// \param DestOffset is an offset from the destination origin measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param DestImgDesc is the destination image descriptor + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels as determined by \p DestImgDesc + void ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent, + ext::oneapi::experimental::image_mem_handle Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + sycl::range<3> CopyExtent); + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. + /// An exception is thrown if either \p Src is incomplete or \p Dest is + /// nullptr. The behavior is undefined if \p Desc is inconsistent with the + /// allocated memory region. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param SrcImgDesc is the source image descriptor + void ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc); + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p Extent are + /// used to determine the sub-region. Pixel size is determined + /// by \p SrcImgDesc + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param SrcOffset is an offset from the origin of source measured in pixels + /// (pixel size determined by \p SrcImgDesc ) + /// \param SrcImgDesc is the source image descriptor + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an offset from the destination origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DestExtent is the extent of the dest memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p SrcImgDesc ) + void + ext_oneapi_copy(ext::oneapi::experimental::image_mem_handle Src, + sycl::range<3> SrcOffset, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, + void *Dest, sycl::range<3> DestOffset, + sycl::range<3> DestExtent, sycl::range<3> CopyExtent); + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. An exception is thrown if either \p Src is nullptr, \p + /// Dest is nullptr, or \p Pitch is inconsistent with hardware requirements. + /// The behavior is undefined if \p Desc is inconsistent with the allocated + /// memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param DeviceImgDesc is the image descriptor (format, order, dimensions). + /// \param DeviceRowPitch is the pitch of the rows on the device. + void ext_oneapi_copy( + void *Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch); + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset , + /// \p DestOffset , and \p Extent are used to determine the sub-region. + /// Pixel size is determined by \p DestImgDesc + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DeviceImgDesc is the device image descriptor + /// \param DeviceRowPitch is the row pitch on the device + /// \param HostExtent is the extent of the dest memory to copy, measured in + /// pixels (pixel size determined by \p DeviceImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p DeviceImgDesc ) + void ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, void *Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, sycl::range<3> HostExtent, + sycl::range<3> CopyExtent); + + /// Instruct the queue with a non-blocking wait on an external semaphore. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + void ext_oneapi_wait_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle + SemaphoreHandle); + + /// Instruct the queue to signal the external semaphore once all previous + /// commands have completed execution. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + void ext_oneapi_signal_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle + SemaphoreHandle); + private: std::shared_ptr MImpl; std::shared_ptr MQueue; diff --git a/sycl/include/sycl/info/aspects.def b/sycl/include/sycl/info/aspects.def index 5c369a72ab743..d1cb99637c0db 100644 --- a/sycl/include/sycl/info/aspects.def +++ b/sycl/include/sycl/info/aspects.def @@ -36,3 +36,14 @@ __SYCL_ASPECT(ext_intel_memory_clock_rate, 38) __SYCL_ASPECT(ext_intel_memory_bus_width, 39) __SYCL_ASPECT(emulated, 40) __SYCL_ASPECT(ext_intel_legacy_image, 41) +__SYCL_ASPECT(ext_oneapi_bindless_images, 42) +__SYCL_ASPECT(ext_oneapi_bindless_images_shared_usm, 43) +__SYCL_ASPECT(ext_oneapi_bindless_images_1d_usm, 44) +__SYCL_ASPECT(ext_oneapi_bindless_images_2d_usm, 45) +__SYCL_ASPECT(ext_oneapi_interop_memory_import, 46) +__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47) +__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48) +__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49) +__SYCL_ASPECT(ext_oneapi_mipmap, 50) +__SYCL_ASPECT(ext_oneapi_mipmap_anisotropy, 51) +__SYCL_ASPECT(ext_oneapi_mipmap_level_reference, 52) diff --git a/sycl/include/sycl/info/ext_oneapi_device_traits.def b/sycl/include/sycl/info/ext_oneapi_device_traits.def index eb0db0ffd6bd5..07e2342cad1e7 100644 --- a/sycl/include/sycl/info/ext_oneapi_device_traits.def +++ b/sycl/include/sycl/info/ext_oneapi_device_traits.def @@ -13,6 +13,25 @@ __SYCL_PARAM_TRAITS_SPEC( ext::oneapi::experimental, device, graph_support, ext::oneapi::experimental::info::graph_support_level, 0 /* No PI device code needed */) + +// Bindless images pitched allocation +__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, + image_row_pitch_align, uint32_t, + PI_EXT_ONEAPI_DEVICE_INFO_IMAGE_PITCH_ALIGN) +__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, + max_image_linear_row_pitch, uint32_t, + PI_EXT_ONEAPI_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH) +__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, + max_image_linear_width, uint32_t, + PI_EXT_ONEAPI_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH) +__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, + max_image_linear_height, uint32_t, + PI_EXT_ONEAPI_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT) + +// Bindles images mipmaps +__SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device, + mipmap_max_anisotropy, float, + PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY) #ifdef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC #undef __SYCL_PARAM_TRAITS_TEMPLATE_SPEC_NEEDS_UNDEF diff --git a/sycl/include/sycl/queue.hpp b/sycl/include/sycl/queue.hpp index 1df6990acf876..9e4370ea70460 100644 --- a/sycl/include/sycl/queue.hpp +++ b/sycl/include/sycl/queue.hpp @@ -1427,6 +1427,683 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase { StartIndex * sizeof(std::remove_all_extents_t)); } + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle wrapper. An exception + /// is thrown if either \p Src is nullptr or \p Dest is incomplete. The + /// behavior is undefined if \p DestImgDesc is inconsistent with the allocated + /// memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a wrapper for an opaque image memory handle to the + /// destination memory. + /// \param DestImgDesc is the image descriptor (format, order, dimensions). + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, ext::oneapi::experimental::image_mem_handle Dest, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. An exception is thrown if either \p + /// Src is nullptr or \p CopyExtent is incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an offset from the origin where the x, y, and z + /// components are measured in bytes, rows, and slices + /// respectively + /// \param SrcExtent is the extent of the source memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param Dest is an opaque image memory handle to the destination memory. + /// \param DestOffset is an offset from the destination origin measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param DestImgDesc is the destination image descriptor + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels as determined by \p DestImgDesc + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent, + ext::oneapi::experimental::image_mem_handle Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + sycl::range<3> CopyExtent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset, + DestImgDesc, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle wrapper. An exception + /// is thrown if either \p Src is nullptr or \p Dest is incomplete. The + /// behavior is undefined if \p DestImgDesc is inconsistent with the allocated + /// memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a wrapper for an opaque image memory handle to the + /// destination memory. + /// \param DestImgDesc is the destination image descriptor + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, ext::oneapi::experimental::image_mem_handle Dest, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. An exception is thrown if either \p + /// Src is nullptr or \p Dest is incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an offset from the origin where the x, y, and z + /// components are measured in bytes, rows, and slices + /// respectively + /// \param SrcExtent is the extent of the source memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param Dest is an opaque image memory handle to the destination memory. + /// \param DestOffset is an offset from the destination origin measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param DestImgDesc is the destination image descriptor + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels as determined by \p DestImgDesc + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent, + ext::oneapi::experimental::image_mem_handle Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + sycl::range<3> CopyExtent, event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset, + DestImgDesc, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle wrapper. An exception + /// is thrown if either \p Src is nullptr or \p Dest is incomplete. The + /// behavior is undefined if \p DestImgDesc is inconsistent with the allocated + /// memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a wrapper for an opaque image memory handle to the + /// destination memory. + /// \param DestImgDesc is the destination image descriptor + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, ext::oneapi::experimental::image_mem_handle Dest, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, Dest, DestImgDesc); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is a USM + /// pointer and \p Dest is an opaque image memory handle. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. An exception is thrown if either \p + /// Src is nullptr or \p Dest is incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an offset from the origin where the x, y, and z + /// components are measured in bytes, rows, and slices + /// respectively + /// \param SrcExtent is the extent of the source memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param Dest is an opaque image memory handle to the destination memory. + /// \param DestOffset is an offset from the destination origin measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param DestImgDesc is the destination image descriptor + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels as determined by \p DestImgDesc + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent, + ext::oneapi::experimental::image_mem_handle Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + sycl::range<3> CopyExtent, const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, SrcOffset, SrcExtent, Dest, DestOffset, + DestImgDesc, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. + /// An exception is thrown if either \p Src is incomplete or \p Dest is + /// nullptr. The behavior is undefined if \p SrcImgDesc is inconsistent with + /// the allocated memory region. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param SrcImgDesc is the source image descriptor. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. Pixel size is determined by \p + /// SrcImgDesc An exception is thrown if either \p Src is nullptr or \p Dest + /// is incomplete. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param SrcOffset is an offset from the origin of source measured in pixels + /// (pixel size determined by \p SrcImgDesc ) + /// \param SrcImgDesc is the source image descriptor + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an offset from the destination origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DestExtent is the extent of the dest memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p SrcImgDesc ) + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest, + sycl::range<3> DestOffset, sycl::range<3> DestExtent, + sycl::range<3> CopyExtent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset, + DestExtent, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. + /// An exception is thrown if either \p Src is incomplete or \p Dest is + /// nullptr. The behavior is undefined if \p SrcImgDesc is inconsistent with + /// the allocated memory region. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param SrcImgDesc is the image descriptor (format, order, dimensions). + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, + event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. Pixel size is determined by \p + /// SrcImgDesc An exception is thrown if either \p Src is nullptr or \p Dest + /// is incomplete. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param SrcOffset is an offset from the origin of source measured in pixels + /// (pixel size determined by \p SrcImgDesc ) + /// \param SrcImgDesc is the source image descriptor + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an offset from the destination origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DestExtent is the extent of the dest memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p SrcImgDesc ) + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest, + sycl::range<3> DestOffset, sycl::range<3> DestExtent, + sycl::range<3> CopyExtent, event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset, + DestExtent, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. + /// An exception is thrown if either \p Src is incomplete or \p Dest is + /// nullptr. The behavior is undefined if \p SrcImgDesc is inconsistent with + /// the allocated memory region. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param SrcImgDesc is the image descriptor (format, order, dimensions). + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, + const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, Dest, SrcImgDesc); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src is an opaque + /// image memory handle and \p Dest is a USM pointer. Allows for a + /// sub-region copy, where \p SrcOffset , \p DestOffset , and \p CopyExtent + /// are used to determine the sub-region. Pixel size is determined by \p + /// SrcImgDesc An exception is thrown if either \p Src is nullptr or \p Dest + /// is incomplete. + /// + /// \param Src is an opaque image memory handle to the source memory. + /// \param SrcOffset is an offset from the origin of source measured in pixels + /// (pixel size determined by \p SrcImgDesc ) + /// \param SrcImgDesc is the source image descriptor + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an offset from the destination origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DestExtent is the extent of the dest memory to copy, measured in + /// pixels (pixel size determined by \p DestImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p SrcImgDesc ) + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest, + sycl::range<3> DestOffset, sycl::range<3> DestExtent, + sycl::range<3> CopyExtent, const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset, + DestExtent, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. An exception is thrown if either \p Src is nullptr, \p + /// Dest is nullptr, or \p Pitch is inconsistent with hardware requirements. + /// The behavior is undefined if \p DeviceImgDesc is inconsistent with the + /// allocated memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param DeviceImgDesc is the image descriptor + /// \param DeviceRowPitch is the DeviceRowPitch of the rows on the device. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset , + /// \p DestOffset , and \p Extent are used to determine the sub-region. + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DeviceImgDesc is the device image descriptor + /// \param DeviceRowPitch is the row pitch on the device + /// \param HostExtent is the extent of the host memory to copy, measured in + /// pixels (pixel size determined by \p DeviceImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p DeviceImgDesc ) + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, void *Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, sycl::range<3> HostExtent, + sycl::range<3> CopyExtent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc, + DeviceRowPitch, HostExtent, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. An exception is thrown if either \p Src is nullptr, \p + /// Dest is nullptr, or \p Pitch is inconsistent with hardware requirements. + /// The behavior is undefined if \p DeviceImgDesc is inconsistent with the + /// allocated memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param DeviceImgDesc is the image descriptor + /// \param DeviceRowPitch is the pitch of the rows on the device. + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset , + /// \p DestOffset , and \p Extent are used to determine the sub-region. + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DeviceImgDesc is the destination image descriptor + /// \param DeviceRowPitch is the row pitch on the device + /// \param HostExtent is the extent of the host memory to copy, measured in + /// pixels (pixel size determined by \p DeviceImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p DeviceImgDesc ) + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, void *Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, sycl::range<3> HostExtent, + sycl::range<3> CopyExtent, event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc, + DeviceRowPitch, HostExtent, CopyExtent); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. An exception is thrown if either \p Src is nullptr, \p + /// Dest is nullptr, or \p Pitch is inconsistent with hardware requirements. + /// The behavior is undefined if \p DeviceImgDesc is inconsistent with the + /// allocated memory region. + /// + /// \param Src is a USM pointer to the source memory. + /// \param Dest is a USM pointer to the destination memory. + /// \param DeviceImgDesc is the image descriptor + /// \param DeviceRowPitch is the pitch of the rows on the device. + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, Dest, DeviceImgDesc, DeviceRowPitch); + }, + CodeLoc); + } + + /// Copies data from one memory region to another, where \p Src and \p Dest + /// are USM pointers. Allows for a sub-region copy, where \p SrcOffset , + /// \p DestOffset , and \p Extent are used to determine the sub-region. + /// An exception is thrown if either \p Src is nullptr or \p Dest is + /// incomplete. + /// + /// \param Src is a USM pointer to the source memory. + /// \param SrcOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param Dest is a USM pointer to the destination memory. + /// \param DestOffset is an destination offset from the origin where the + /// x, y, and z components are measured in bytes, rows, + /// and slices respectively + /// \param DeviceImgDesc is the destination image descriptor + /// \param DeviceRowPitch is the row pitch on the device + /// \param HostExtent is the extent of the host memory to copy, measured in + /// pixels (pixel size determined by \p DeviceImgDesc ) + /// \param CopyExtent is the width, height, and depth of the region to copy + /// measured in pixels (pixel size determined by + /// \p DeviceImgDesc ) + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the copy operation. + event ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, void *Dest, + sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, sycl::range<3> HostExtent, + sycl::range<3> CopyExtent, const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_copy(Src, SrcOffset, Dest, DestOffset, DeviceImgDesc, + DeviceRowPitch, HostExtent, CopyExtent); + }, + CodeLoc); + } + + /// Instruct the queue with a non-blocking wait on an external semaphore. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \return an event representing the wait operation. + event ext_oneapi_wait_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + + /// Instruct the queue with a non-blocking wait on an external semaphore. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the wait operation. + event ext_oneapi_wait_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + + /// Instruct the queue with a non-blocking wait on an external semaphore. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the wait operation. + event ext_oneapi_wait_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_wait_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + + /// Instruct the queue to signal the external semaphore once all previous + /// commands have completed execution. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \return an event representing the signal operation. + event ext_oneapi_signal_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + + /// Instruct the queue to signal the external semaphore once all previous + /// commands have completed execution. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \param DepEvent is an event that specifies the kernel dependencies. + /// \return an event representing the signal operation. + event ext_oneapi_signal_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + event DepEvent, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvent); + CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + + /// Instruct the queue to signal the external semaphore once all previous + /// commands have completed execution. + /// An exception is thrown if \p SemaphoreHandle is incomplete. + /// + /// \param SemaphoreHandle is an opaque external interop semaphore handle + /// \param DepEvents is a vector of events that specifies the kernel + /// dependencies. + /// \return an event representing the signal operation. + event ext_oneapi_signal_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle, + const std::vector &DepEvents, + const detail::code_location &CodeLoc = detail::code_location::current()) { + detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc); + return submit( + [&](handler &CGH) { + CGH.depends_on(DepEvents); + CGH.ext_oneapi_signal_external_semaphore(SemaphoreHandle); + }, + CodeLoc); + } + /// single_task version with a kernel represented as a lambda. /// /// \param Properties is the kernel properties. diff --git a/sycl/include/sycl/sycl.hpp b/sycl/include/sycl/sycl.hpp index f51d9107e588b..3563e19128429 100644 --- a/sycl/include/sycl/sycl.hpp +++ b/sycl/include/sycl/sycl.hpp @@ -64,6 +64,7 @@ #include #include #include +#include #include #include #include diff --git a/sycl/source/CMakeLists.txt b/sycl/source/CMakeLists.txt index 79f4535ae0175..4e837d551f1c3 100644 --- a/sycl/source/CMakeLists.txt +++ b/sycl/source/CMakeLists.txt @@ -153,6 +153,7 @@ set(SYCL_SOURCES "backend.cpp" "detail/accessor_impl.cpp" "detail/allowlist.cpp" + "detail/bindless_images.cpp" "detail/buffer_impl.cpp" "detail/builtins_common.cpp" "detail/builtins_geometric.cpp" diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp new file mode 100644 index 0000000000000..92b50c9d15fba --- /dev/null +++ b/sycl/source/detail/bindless_images.cpp @@ -0,0 +1,750 @@ +//==----------- bindless_images.hpp --- SYCL bindless images ---------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +#include + +#include +#include +#include +#include + +#include + +namespace sycl { +inline namespace _V1 { +namespace ext::oneapi::experimental { + +void populate_pi_structs(const image_descriptor &desc, pi_image_desc &piDesc, + pi_image_format &piFormat, size_t pitch = 0) { + piDesc = {}; + piDesc.image_width = desc.width; + piDesc.image_height = desc.height; + piDesc.image_depth = desc.depth; + piDesc.image_type = desc.depth > 0 ? PI_MEM_TYPE_IMAGE3D + : (desc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + piDesc.image_row_pitch = pitch; + piDesc.image_array_size = 0; + piDesc.image_slice_pitch = 0; + piDesc.num_mip_levels = desc.num_levels; + piDesc.num_samples = 0; + piDesc.buffer = nullptr; + + piFormat = {}; + piFormat.image_channel_data_type = + sycl::detail::convertChannelType(desc.channel_type); + piFormat.image_channel_order = + sycl::detail::convertChannelOrder(desc.channel_order); +} + +detail::image_mem_impl::image_mem_impl(const image_descriptor &desc, + const device &syclDevice, + const context &syclContext) + : descriptor(desc), syclDevice(syclDevice), syclContext(syclContext) { + if (desc.type == image_type::mipmap) { + handle = alloc_mipmap_mem(desc, syclDevice, syclContext); + } else { + handle = alloc_image_mem(desc, syclDevice, syclContext); + } +} + +detail::image_mem_impl::~image_mem_impl() { + if (handle.raw_handle != nullptr) { + if (descriptor.type == image_type::mipmap) { + free_mipmap_mem(handle, syclDevice, syclContext); + } else { + free_image_mem(handle, syclDevice, syclContext); + } + } +} + +__SYCL_EXPORT +image_mem::image_mem(const image_descriptor &desc, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + impl = + std::make_shared(desc, syclDevice, syclContext); +} + +__SYCL_EXPORT +image_mem::image_mem(const image_descriptor &desc, const sycl::queue &syclQueue) + : image_mem(desc, syclQueue.get_device(), syclQueue.get_context()) {} + +__SYCL_EXPORT sycl::range<3> image_mem::get_range() const { + auto desc = impl->get_descriptor(); + return {desc.width, desc.height, desc.depth}; +} + +__SYCL_EXPORT sycl::image_channel_type image_mem::get_channel_type() const { + return impl->get_descriptor().channel_type; +} + +__SYCL_EXPORT sycl::image_channel_order image_mem::get_channel_order() const { + return impl->get_descriptor().channel_order; +} + +__SYCL_EXPORT unsigned int image_mem::get_num_channels() const { + return sycl::detail::getImageNumberChannels( + impl->get_descriptor().channel_order); +} + +__SYCL_EXPORT image_type image_mem::get_type() const { + return impl->get_descriptor().type; +} + +__SYCL_EXPORT image_mem_handle +image_mem::get_mip_level_mem_handle(const unsigned int level) const { + return ext::oneapi::experimental::get_mip_level_mem_handle( + impl->get_handle(), level, impl->get_device(), impl->get_context()); +} + +__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_image_handle piImageHandle = imageHandle.raw_handle; + + Plugin->call( + C, Device, piImageHandle); +} + +__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle, + const sycl::queue &syclQueue) { + destroy_image_handle(imageHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_image_handle piImageHandle = imageHandle.raw_handle; + + Plugin->call( + C, Device, piImageHandle); +} + +__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle, + const sycl::queue &syclQueue) { + destroy_image_handle(imageHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT image_mem_handle +alloc_image_mem(const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + // Non-mipmap images must have only 1 level + if (desc.num_levels != 1) + throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), + "Image number of levels must be 1"); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat); + + image_mem_handle retHandle; + + // Call impl. + Plugin->call( + C, Device, &piFormat, &piDesc, &retHandle.raw_handle); + + return retHandle; +} + +__SYCL_EXPORT image_mem_handle alloc_image_mem(const image_descriptor &desc, + const sycl::queue &syclQueue) { + return alloc_image_mem(desc, syclQueue.get_device(), syclQueue.get_context()); +} + +__SYCL_EXPORT image_mem_handle +alloc_mipmap_mem(const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + // Mipmaps must have more than one level + if (desc.num_levels <= 1) + throw sycl::exception(sycl::make_error_code(sycl::errc::invalid), + "Mipmap number of levels must be 2 or more"); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat); + + // Call impl. + image_mem_handle retHandle; + Plugin->call( + C, Device, &piFormat, &piDesc, &retHandle.raw_handle); + + return retHandle; +} + +__SYCL_EXPORT image_mem_handle alloc_mipmap_mem(const image_descriptor &desc, + const sycl::queue &syclQueue) { + return alloc_mipmap_mem(desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT image_mem_handle get_mip_level_mem_handle( + const image_mem_handle mipMem, unsigned int level, + const sycl::device &syclDevice, const sycl::context &syclContext) { + + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + // Call impl. + image_mem_handle individual_image; + Plugin->call( + C, Device, mipMem.raw_handle, level, &individual_image.raw_handle); + + return individual_image; +} + +__SYCL_EXPORT image_mem_handle +get_mip_level_mem_handle(const image_mem_handle mipMem, unsigned int level, + const sycl::queue &syclQueue) { + return get_mip_level_mem_handle(mipMem, level, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void free_image_mem(image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + Plugin->call( + C, Device, memHandle.raw_handle); +} + +__SYCL_EXPORT void free_image_mem(image_mem_handle memHandle, + const sycl::queue &syclQueue) { + free_image_mem(memHandle, syclQueue.get_device(), syclQueue.get_context()); +} + +__SYCL_EXPORT void free_mipmap_mem(image_mem_handle memoryHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + Plugin->call( + C, Device, memoryHandle.raw_handle); +} + +__SYCL_EXPORT void free_mipmap_mem(image_mem_handle memoryHandle, + const sycl::queue &syclQueue) { + free_mipmap_mem(memoryHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem &imgMem, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext) { + return create_image(imgMem.get_handle(), desc, syclDevice, syclContext); +} + +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem &imgMem, const image_descriptor &desc, + const sycl::queue &syclQueue) { + return create_image(imgMem.get_handle(), desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem_handle memHandle, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat); + + // Call impl. + pi_image_handle piImageHandle; + pi_mem piImage; + Plugin->call( + C, Device, memHandle.raw_handle, &piFormat, &piDesc, &piImage, + &piImageHandle); + + return unsampled_image_handle{piImageHandle}; +} + +__SYCL_EXPORT unsampled_image_handle +create_image(image_mem_handle memHandle, const image_descriptor &desc, + const sycl::queue &syclQueue) { + return create_image(memHandle, desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT sampled_image_handle +create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + const pi_sampler_properties sProps[] = { + PI_SAMPLER_INFO_NORMALIZED_COORDS, + static_cast(sampler.coordinate), + PI_SAMPLER_INFO_ADDRESSING_MODE, + static_cast(sampler.addressing), + PI_SAMPLER_INFO_FILTER_MODE, + static_cast(sampler.filtering), + PI_SAMPLER_INFO_MIP_FILTER_MODE, + static_cast(sampler.mipmap_filtering), + 0}; + + pi_sampler piSampler = {}; + Plugin->call( + C, sProps, sampler.min_mipmap_level_clamp, sampler.max_mipmap_level_clamp, + sampler.max_anisotropy, &piSampler); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat); + + // Call impl. + pi_image_handle piImageHandle; + pi_mem piImage; + Plugin->call( + C, Device, memHandle.raw_handle, &piFormat, &piDesc, piSampler, &piImage, + &piImageHandle); + + return sampled_image_handle{piImageHandle}; +} + +__SYCL_EXPORT sampled_image_handle +create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue) { + return create_image(memHandle, sampler, desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT sampled_image_handle +create_image(image_mem &imgMem, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext) { + return create_image(imgMem.get_handle().raw_handle, 0 /*pitch*/, sampler, + desc, syclDevice, syclContext); +} + +__SYCL_EXPORT sampled_image_handle +create_image(image_mem &imgMem, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue) { + return create_image(imgMem.get_handle().raw_handle, 0 /*pitch*/, sampler, + desc, syclQueue.get_device(), syclQueue.get_context()); +} + +__SYCL_EXPORT sampled_image_handle +create_image(void *devPtr, size_t pitch, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::device &syclDevice, + const sycl::context &syclContext) { + + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + const pi_sampler_properties sProps[] = { + PI_SAMPLER_INFO_NORMALIZED_COORDS, + static_cast(sampler.coordinate), + PI_SAMPLER_INFO_ADDRESSING_MODE, + static_cast(sampler.addressing), + PI_SAMPLER_INFO_FILTER_MODE, + static_cast(sampler.filtering), + PI_SAMPLER_INFO_MIP_FILTER_MODE, + static_cast(sampler.mipmap_filtering), + 0}; + + pi_sampler piSampler = {}; + Plugin->call( + C, sProps, sampler.min_mipmap_level_clamp, sampler.max_mipmap_level_clamp, + sampler.max_anisotropy, &piSampler); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat, pitch); + + // Call impl. + pi_mem piImage; + pi_image_handle piImageHandle; + Plugin->call( + C, Device, devPtr, &piFormat, &piDesc, piSampler, &piImage, + &piImageHandle); + + return sampled_image_handle{piImageHandle}; +} + +__SYCL_EXPORT sampled_image_handle +create_image(void *devPtr, size_t pitch, const bindless_image_sampler &sampler, + const image_descriptor &desc, const sycl::queue &syclQueue) { + return create_image(devPtr, pitch, sampler, desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +template <> +__SYCL_EXPORT interop_mem_handle import_external_memory( + external_mem_descriptor externalMem, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + pi_interop_mem_handle piInteropMem; + Plugin->call( + C, Device, externalMem.size_in_bytes, + externalMem.external_handle.file_descriptor, &piInteropMem); + + return interop_mem_handle{piInteropMem}; +} + +template <> +__SYCL_EXPORT interop_mem_handle import_external_memory( + external_mem_descriptor externalMem, + const sycl::queue &syclQueue) { + return import_external_memory( + externalMem, syclQueue.get_device(), syclQueue.get_context()); +} + +__SYCL_EXPORT image_mem_handle map_external_memory_array( + interop_mem_handle memHandle, const image_descriptor &desc, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + pi_image_desc piDesc; + pi_image_format piFormat; + populate_pi_structs(desc, piDesc, piFormat); + + pi_interop_mem_handle piInteropMem{memHandle.raw_handle}; + + image_mem_handle retHandle; + Plugin->call( + C, Device, &piFormat, &piDesc, piInteropMem, &retHandle.raw_handle); + + return image_mem_handle{retHandle}; +} + +__SYCL_EXPORT image_mem_handle map_external_memory_array( + interop_mem_handle memHandle, const image_descriptor &desc, + const sycl::queue &syclQueue) { + return map_external_memory_array(memHandle, desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void release_external_memory(interop_mem_handle interopMem, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + Plugin->call( + C, Device, (pi_interop_mem_handle)interopMem.raw_handle); +} + +__SYCL_EXPORT void release_external_memory(interop_mem_handle interopMem, + const sycl::queue &syclQueue) { + release_external_memory(interopMem, syclQueue.get_device(), + syclQueue.get_context()); +} + +template <> +__SYCL_EXPORT interop_semaphore_handle import_external_semaphore( + external_semaphore_descriptor externalSemaphoreDesc, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + + pi_interop_semaphore_handle piInteropSemaphore; + + Plugin->call( + C, Device, externalSemaphoreDesc.external_handle.file_descriptor, + &piInteropSemaphore); + + return interop_semaphore_handle{piInteropSemaphore}; +} + +template <> +__SYCL_EXPORT interop_semaphore_handle import_external_semaphore( + external_semaphore_descriptor externalSemaphoreDesc, + const sycl::queue &syclQueue) { + return import_external_semaphore( + externalSemaphoreDesc, syclQueue.get_device(), syclQueue.get_context()); +} + +__SYCL_EXPORT void +destroy_external_semaphore(interop_semaphore_handle semaphoreHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_context C = CtxImpl->getHandleRef(); + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + pi_device Device = DevImpl->getHandleRef(); + + Plugin->call( + C, Device, (pi_interop_semaphore_handle)semaphoreHandle.raw_handle); +} + +__SYCL_EXPORT void +destroy_external_semaphore(interop_semaphore_handle semaphoreHandle, + const sycl::queue &syclQueue) { + destroy_external_semaphore(semaphoreHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::ignore = syclDevice; + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + size_t Width, Height, Depth; + + Plugin->call( + memHandle.raw_handle, PI_IMAGE_INFO_WIDTH, &Width, nullptr); + + Plugin->call( + memHandle.raw_handle, PI_IMAGE_INFO_HEIGHT, &Height, nullptr); + + Plugin->call( + memHandle.raw_handle, PI_IMAGE_INFO_DEPTH, &Depth, nullptr); + + return {Width, Height, Depth}; +} + +__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle, + const sycl::queue &syclQueue) { + return get_image_range(memHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT sycl::image_channel_type +get_image_channel_type(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::ignore = syclDevice; + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + + pi_image_format PIFormat; + + Plugin->call( + memHandle.raw_handle, PI_IMAGE_INFO_FORMAT, &PIFormat, nullptr); + + image_channel_type ChannelType = + sycl::detail::convertChannelType(PIFormat.image_channel_data_type); + + return ChannelType; +} + +__SYCL_EXPORT sycl::image_channel_type +get_image_channel_type(const image_mem_handle memHandle, + const sycl::queue &syclQueue) { + return get_image_channel_type(memHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + size_t widthInBytes, size_t height, + unsigned int elementSizeBytes, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + void *RetVal = nullptr; + if (widthInBytes == 0 || height == 0 || elementSizeBytes == 0) { + throw sycl::exception(sycl::make_error_code(sycl::errc::memory_allocation), + "Cannot allocate pitched memory with zero size!"); + } + + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + if (CtxImpl->is_host()) { + throw sycl::exception(sycl::make_error_code(sycl::errc::memory_allocation), + "Cannot allocate pitched memory on host!"); + } + + pi_context PiContext = CtxImpl->getHandleRef(); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_device PiDevice; + + PiDevice = sycl::detail::getSyclObjImpl(syclDevice)->getHandleRef(); + + Plugin->call( + &RetVal, resultPitch, PiContext, PiDevice, nullptr, widthInBytes, height, + elementSizeBytes); + + return RetVal; +} + +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + size_t widthInBytes, size_t height, + unsigned int elementSizeBytes, + const sycl::queue &syclQueue) { + return pitched_alloc_device(resultPitch, widthInBytes, height, + elementSizeBytes, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + const image_descriptor &desc, + const sycl::queue &syclQueue) { + return pitched_alloc_device(resultPitch, desc, syclQueue.get_device(), + syclQueue.get_context()); +} + +__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch, + const image_descriptor &desc, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + uint8_t numChannels = + sycl::detail::getImageNumberChannels(desc.channel_order); + unsigned int elementSizeBytes = + sycl::detail::getImageElementSize(numChannels, desc.channel_type); + + size_t widthInBytes = desc.width * elementSizeBytes; + size_t height = desc.height; + + return pitched_alloc_device(resultPitch, widthInBytes, height, + elementSizeBytes, syclDevice, syclContext); +} + +__SYCL_EXPORT unsigned int +get_image_num_channels(const image_mem_handle memHandle, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::ignore = syclDevice; + + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::PluginPtr &Plugin = CtxImpl->getPlugin(); + pi_image_format PIFormat; + + Plugin->call( + memHandle.raw_handle, PI_IMAGE_INFO_FORMAT, &PIFormat, nullptr); + + image_channel_order Order = + sycl::detail::convertChannelOrder(PIFormat.image_channel_order); + + return static_cast(sycl::detail::getImageNumberChannels(Order)); +} + +__SYCL_EXPORT unsigned int +get_image_num_channels(const image_mem_handle memHandle, + const sycl::queue &syclQueue) { + return get_image_num_channels(memHandle, syclQueue.get_device(), + syclQueue.get_context()); +} + +} // namespace ext::oneapi::experimental +} // namespace _V1 +} // namespace sycl diff --git a/sycl/source/detail/device_impl.cpp b/sycl/source/detail/device_impl.cpp index 5c1e854bff183..fa902b038db7e 100644 --- a/sycl/source/detail/device_impl.cpp +++ b/sycl/source/detail/device_impl.cpp @@ -457,6 +457,95 @@ bool device_impl::has(aspect Aspect) const { &legacy_image_support, nullptr) == PI_SUCCESS; return call_successful && legacy_image_support; } + case aspect::ext_oneapi_bindless_images: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_bindless_images_shared_usm: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, + PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_SHARED_USM_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_bindless_images_1d_usm: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_1D_USM_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_bindless_images_2d_usm: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_IMAGES_2D_USM_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_interop_memory_import: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_MEMORY_IMPORT_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_interop_memory_export: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_MEMORY_EXPORT_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_interop_semaphore_import: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_SEMAPHORE_IMPORT_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_interop_semaphore_export: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_SEMAPHORE_EXPORT_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_mipmap: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_SUPPORT, sizeof(pi_bool), + &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_mipmap_anisotropy: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_ANISOTROPY_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } + case aspect::ext_oneapi_mipmap_level_reference: { + pi_bool support = PI_FALSE; + bool call_successful = + getPlugin()->call_nocheck( + MDevice, PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_LEVEL_REFERENCE_SUPPORT, + sizeof(pi_bool), &support, nullptr) == PI_SUCCESS; + return call_successful && support; + } } throw runtime_error("This device aspect has not been implemented yet.", PI_ERROR_INVALID_DEVICE); diff --git a/sycl/source/detail/device_info.hpp b/sycl/source/detail/device_info.hpp index 0f086248269f2..180fd6f6933f4 100644 --- a/sycl/source/detail/device_info.hpp +++ b/sycl/source/detail/device_info.hpp @@ -1821,6 +1821,45 @@ get_device_info_host() { return ext::oneapi::experimental::info::graph_support_level::unsupported; } +template <> +inline uint32_t get_device_info_host< + ext::oneapi::experimental::info::device::image_row_pitch_align>() { + throw runtime_error("Obtaining image pitch alignment is not " + "supported on HOST device", + PI_ERROR_INVALID_DEVICE); +} + +template <> +inline uint32_t get_device_info_host< + ext::oneapi::experimental::info::device::max_image_linear_row_pitch>() { + throw runtime_error("Obtaining max image linear pitch is not " + "supported on HOST device", + PI_ERROR_INVALID_DEVICE); +} + +template <> +inline uint32_t get_device_info_host< + ext::oneapi::experimental::info::device::max_image_linear_width>() { + throw runtime_error("Obtaining max image linear width is not " + "supported on HOST device", + PI_ERROR_INVALID_DEVICE); +} + +template <> +inline uint32_t get_device_info_host< + ext::oneapi::experimental::info::device::max_image_linear_height>() { + throw runtime_error("Obtaining max image linear height is not " + "supported on HOST device", + PI_ERROR_INVALID_DEVICE); +} + +template <> +inline float get_device_info_host< + ext::oneapi::experimental::info::device::mipmap_max_anisotropy>() { + throw runtime_error("Bindless image mipaps are not supported on HOST device", + PI_ERROR_INVALID_DEVICE); +} + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/graph_impl.hpp b/sycl/source/detail/graph_impl.hpp index 1972d19b261ad..5526aeaccef44 100644 --- a/sycl/source/detail/graph_impl.hpp +++ b/sycl/source/detail/graph_impl.hpp @@ -149,6 +149,12 @@ class node_impl { return createCGCopy(); case sycl::detail::CG::ReadWriteHostPipe: return createCGCopy(); + case sycl::detail::CG::CopyImage: + return createCGCopy(); + case sycl::detail::CG::SemaphoreSignal: + return createCGCopy(); + case sycl::detail::CG::SemaphoreWait: + return createCGCopy(); case sycl::detail::CG::ExecCommandBuffer: assert(false && "Error: Command graph submission should not be a node in a graph"); diff --git a/sycl/source/detail/handler_impl.hpp b/sycl/source/detail/handler_impl.hpp index 0efca386b17d3..d98602ab02e35 100644 --- a/sycl/source/detail/handler_impl.hpp +++ b/sycl/source/detail/handler_impl.hpp @@ -104,6 +104,19 @@ class handler_impl { sycl::detail::pi::PiKernelCacheConfig MKernelCacheConfig = PI_EXT_KERNEL_EXEC_INFO_CACHE_DEFAULT; + + // Extra information for bindless image copy + sycl::detail::pi::PiMemImageDesc MImageDesc; + sycl::detail::pi::PiMemImageFormat MImageFormat; + sycl::detail::pi::PiImageCopyFlags MImageCopyFlags; + + sycl::detail::pi::PiImageOffset MSrcOffset; + sycl::detail::pi::PiImageOffset MDestOffset; + sycl::detail::pi::PiImageRegion MHostExtent; + sycl::detail::pi::PiImageRegion MCopyExtent; + + // Extra information for semaphore interoperability + sycl::detail::pi::PiInteropSemaphoreHandle MInteropSemaphoreHandle; }; } // namespace detail diff --git a/sycl/source/detail/memory_manager.cpp b/sycl/source/detail/memory_manager.cpp index 9591f08a5f568..c55591e59685f 100644 --- a/sycl/source/detail/memory_manager.cpp +++ b/sycl/source/detail/memory_manager.cpp @@ -1432,6 +1432,37 @@ void MemoryManager::ext_oneapi_copy_usm_cmd_buffer( OutSyncPoint); } +void MemoryManager::copy_image_bindless( + void *Src, QueueImplPtr Queue, void *Dst, + const sycl::detail::pi::PiMemImageDesc &Desc, + const sycl::detail::pi::PiMemImageFormat &Format, + const sycl::detail::pi::PiImageCopyFlags Flags, + sycl::detail::pi::PiImageOffset SrcOffset, + sycl::detail::pi::PiImageOffset DstOffset, + sycl::detail::pi::PiImageRegion HostExtent, + sycl::detail::pi::PiImageRegion CopyExtent, + const std::vector &DepEvents, + sycl::detail::pi::PiEvent *OutEvent) { + + assert(!Queue->getContextImplPtr()->is_host() && + "Host queue not supported in copy_image_bindless."); + assert((Flags == (sycl::detail::pi::PiImageCopyFlags) + ext::oneapi::experimental::image_copy_flags::HtoD || + Flags == (sycl::detail::pi::PiImageCopyFlags) + ext::oneapi::experimental::image_copy_flags::DtoH) && + "Invalid flags passed to copy_image_bindless."); + if (!Dst || !Src) + throw sycl::exception( + sycl::make_error_code(errc::invalid), + "NULL pointer argument in bindless image copy operation."); + + const detail::PluginPtr &Plugin = Queue->getPlugin(); + Plugin->call( + Queue->getHandleRef(), Dst, Src, &Format, &Desc, Flags, &SrcOffset, + &DstOffset, &CopyExtent, &HostExtent, DepEvents.size(), DepEvents.data(), + OutEvent); +} + } // namespace detail } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/memory_manager.hpp b/sycl/source/detail/memory_manager.hpp index a824e70deee3e..27369b2f636e6 100644 --- a/sycl/source/detail/memory_manager.hpp +++ b/sycl/source/detail/memory_manager.hpp @@ -227,6 +227,18 @@ class __SYCL_EXPORT MemoryManager { sycl::detail::pi::PiExtCommandBuffer CommandBuffer, size_t Len, void *DstMem, std::vector Deps, sycl::detail::pi::PiExtSyncPoint *OutSyncPoint); + + static void + copy_image_bindless(void *Src, QueueImplPtr Queue, void *Dst, + const sycl::detail::pi::PiMemImageDesc &Desc, + const sycl::detail::pi::PiMemImageFormat &Format, + const sycl::detail::pi::PiImageCopyFlags Flags, + sycl::detail::pi::PiImageOffset SrcOffset, + sycl::detail::pi::PiImageOffset DstOffset, + sycl::detail::pi::PiImageRegion CopyExtent, + sycl::detail::pi::PiImageRegion HostExtent, + const std::vector &DepEvents, + sycl::detail::pi::PiEvent *OutEvent); }; } // namespace detail } // namespace _V1 diff --git a/sycl/source/detail/scheduler/commands.cpp b/sycl/source/detail/scheduler/commands.cpp index bb3c7072c0157..f34b8f18acfe1 100644 --- a/sycl/source/detail/scheduler/commands.cpp +++ b/sycl/source/detail/scheduler/commands.cpp @@ -3037,6 +3037,46 @@ pi_int32 ExecCGCommand::enqueueImpQueue() { RawEvents.size(), RawEvents.empty() ? nullptr : &RawEvents[0], Event); } + case CG::CGTYPE::CopyImage: { + CGCopyImage *Copy = (CGCopyImage *)MCommandGroup.get(); + + sycl::detail::pi::PiMemImageDesc Desc = Copy->getDesc(); + + MemoryManager::copy_image_bindless( + Copy->getSrc(), MQueue, Copy->getDst(), Desc, Copy->getFormat(), + Copy->getCopyFlags(), Copy->getSrcOffset(), Copy->getDstOffset(), + Copy->getHostExtent(), Copy->getCopyExtent(), std::move(RawEvents), + Event); + return PI_SUCCESS; + } + case CG::CGTYPE::SemaphoreWait: { + CGSemaphoreWait *SemWait = (CGSemaphoreWait *)MCommandGroup.get(); + if (MQueue->getDeviceImplPtr()->is_host()) { + // NOP for host device. + return PI_SUCCESS; + } + + const detail::PluginPtr &Plugin = MQueue->getPlugin(); + Plugin->call( + MQueue->getHandleRef(), SemWait->getInteropSemaphoreHandle(), 0, + nullptr, nullptr); + + return PI_SUCCESS; + } + case CG::CGTYPE::SemaphoreSignal: { + CGSemaphoreSignal *SemSignal = (CGSemaphoreSignal *)MCommandGroup.get(); + if (MQueue->getDeviceImplPtr()->is_host()) { + // NOP for host device. + return PI_SUCCESS; + } + + const detail::PluginPtr &Plugin = MQueue->getPlugin(); + Plugin->call( + MQueue->getHandleRef(), SemSignal->getInteropSemaphoreHandle(), 0, + nullptr, nullptr); + + return PI_SUCCESS; + } case CG::CGTYPE::None: throw sycl::exception(sycl::make_error_code(sycl::errc::runtime), "CG type not implemented. " + diff --git a/sycl/source/feature_test.hpp.in b/sycl/source/feature_test.hpp.in index bd2371651e4ab..116fa2312e9e3 100644 --- a/sycl/source/feature_test.hpp.in +++ b/sycl/source/feature_test.hpp.in @@ -40,6 +40,7 @@ inline namespace _V1 { #define SYCL_EXT_ONEAPI_QUEUE_PRIORITY 1 #define SYCL_EXT_ONEAPI_ENQUEUE_BARRIER 1 #define SYCL_EXT_ONEAPI_FREE_FUNCTION_QUERIES 1 +#define SYCL_EXT_ONEAPI_BINDLESS_IMAGES 1 #define SYCL_EXT_ONEAPI_GROUP_ALGORITHMS 1 #define SYCL_EXT_ONEAPI_GROUP_SORT 1 #define SYCL_EXT_ONEAPI_MAX_WORK_GROUP_QUERY 1 diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 34075982c4948..95db5c5eb66af 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -12,6 +12,7 @@ #include #include #include +#include #include #include #include @@ -40,6 +41,34 @@ bool isDeviceGlobalUsedInKernel(const void *DeviceGlobalPtr) { return DGEntry && !DGEntry->MImageIdentifiers.empty(); } +sycl::detail::pi::PiImageCopyFlags +getPiImageCopyFlags(sycl::usm::alloc SrcPtrType, sycl::usm::alloc DstPtrType) { + if (DstPtrType == sycl::usm::alloc::device) { + // Dest is on device + if (SrcPtrType == sycl::usm::alloc::device) + return sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_DEVICE_TO_DEVICE; + if (SrcPtrType == sycl::usm::alloc::host || + SrcPtrType == sycl::usm::alloc::unknown) + return sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_HOST_TO_DEVICE; + throw sycl::exception(make_error_code(errc::invalid), + "Unknown copy source location"); + } + if (DstPtrType == sycl::usm::alloc::host || + DstPtrType == sycl::usm::alloc::unknown) { + // Dest is on host + if (SrcPtrType == sycl::usm::alloc::device) + return sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_DEVICE_TO_HOST; + if (SrcPtrType == sycl::usm::alloc::host || + SrcPtrType == sycl::usm::alloc::unknown) + throw sycl::exception(make_error_code(errc::invalid), + "Cannot copy image from host to host"); + throw sycl::exception(make_error_code(errc::invalid), + "Unknown copy source location"); + } + throw sycl::exception(make_error_code(errc::invalid), + "Unknown copy destination location"); +} + } // namespace detail handler::handler(std::shared_ptr Queue, bool IsHost) @@ -373,6 +402,20 @@ event handler::finalize() { return MLastEvent; } break; + case detail::CG::CopyImage: + CommandGroup.reset(new detail::CGCopyImage( + MSrcPtr, MDstPtr, MImpl->MImageDesc, MImpl->MImageFormat, + MImpl->MImageCopyFlags, MImpl->MSrcOffset, MImpl->MDestOffset, + MImpl->MHostExtent, MImpl->MCopyExtent, std::move(CGData), MCodeLoc)); + break; + case detail::CG::SemaphoreWait: + CommandGroup.reset(new detail::CGSemaphoreWait( + MImpl->MInteropSemaphoreHandle, std::move(CGData), MCodeLoc)); + break; + case detail::CG::SemaphoreSignal: + CommandGroup.reset(new detail::CGSemaphoreSignal( + MImpl->MInteropSemaphoreHandle, std::move(CGData), MCodeLoc)); + break; case detail::CG::None: if (detail::pi::trace(detail::pi::TraceLevel::PI_TRACE_ALL)) { std::cout << "WARNING: An empty command group is submitted." << std::endl; @@ -842,6 +885,220 @@ void handler::ext_oneapi_memset2d_impl(void *Dest, size_t DestPitch, int Value, setType(detail::CG::Memset2DUSM); } +void handler::ext_oneapi_copy( + void *Src, ext::oneapi::experimental::image_mem_handle Dest, + const ext::oneapi::experimental::image_descriptor &Desc) { + MSrcPtr = Src; + MDstPtr = Dest.raw_handle; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = Desc.width; + PiDesc.image_height = Desc.height; + PiDesc.image_depth = Desc.depth; + PiDesc.image_type = Desc.depth > 0 ? PI_MEM_TYPE_IMAGE3D + : (Desc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(Desc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(Desc.channel_order); + + MImpl->MSrcOffset = {0, 0, 0}; + MImpl->MDestOffset = {0, 0, 0}; + MImpl->MCopyExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MHostExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = + sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_HOST_TO_DEVICE; + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, sycl::range<3> SrcExtent, + ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DestImgDesc, + sycl::range<3> CopyExtent) { + + MSrcPtr = Src; + MDstPtr = Dest.raw_handle; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = DestImgDesc.width; + PiDesc.image_height = DestImgDesc.height; + PiDesc.image_depth = DestImgDesc.depth; + PiDesc.image_type = DestImgDesc.depth > 0 + ? PI_MEM_TYPE_IMAGE3D + : (DestImgDesc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(DestImgDesc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(DestImgDesc.channel_order); + + MImpl->MSrcOffset = {SrcOffset[0], SrcOffset[1], SrcOffset[2]}; + MImpl->MDestOffset = {DestOffset[0], DestOffset[1], DestOffset[2]}; + MImpl->MCopyExtent = {CopyExtent[0], CopyExtent[1], CopyExtent[2]}; + MImpl->MHostExtent = {SrcExtent[0], SrcExtent[1], SrcExtent[2]}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = + sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_HOST_TO_DEVICE; + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &Desc) { + MSrcPtr = Src.raw_handle; + MDstPtr = Dest; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = Desc.width; + PiDesc.image_height = Desc.height; + PiDesc.image_depth = Desc.depth; + PiDesc.image_type = Desc.depth > 0 ? PI_MEM_TYPE_IMAGE3D + : (Desc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(Desc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(Desc.channel_order); + + MImpl->MSrcOffset = {0, 0, 0}; + MImpl->MDestOffset = {0, 0, 0}; + MImpl->MCopyExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MHostExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = + sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_DEVICE_TO_HOST; + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_copy( + ext::oneapi::experimental::image_mem_handle Src, sycl::range<3> SrcOffset, + const ext::oneapi::experimental::image_descriptor &SrcImgDesc, void *Dest, + sycl::range<3> DestOffset, sycl::range<3> DestExtent, + sycl::range<3> CopyExtent) { + MSrcPtr = Src.raw_handle; + MDstPtr = Dest; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = SrcImgDesc.width; + PiDesc.image_height = SrcImgDesc.height; + PiDesc.image_depth = SrcImgDesc.depth; + PiDesc.image_type = + SrcImgDesc.depth > 0 + ? PI_MEM_TYPE_IMAGE3D + : (SrcImgDesc.height > 0 ? PI_MEM_TYPE_IMAGE2D : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(SrcImgDesc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(SrcImgDesc.channel_order); + + MImpl->MSrcOffset = {SrcOffset[0], SrcOffset[1], SrcOffset[2]}; + MImpl->MDestOffset = {DestOffset[0], DestOffset[1], DestOffset[2]}; + MImpl->MCopyExtent = {CopyExtent[0], CopyExtent[1], CopyExtent[2]}; + MImpl->MHostExtent = {DestExtent[0], DestExtent[1], DestExtent[2]}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = + sycl::detail::pi::PiImageCopyFlags::PI_IMAGE_COPY_DEVICE_TO_HOST; + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_copy( + void *Src, void *Dest, + const ext::oneapi::experimental::image_descriptor &Desc, size_t Pitch) { + MSrcPtr = Src; + MDstPtr = Dest; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = Desc.width; + PiDesc.image_height = Desc.height; + PiDesc.image_depth = Desc.depth; + PiDesc.image_type = Desc.depth > 0 ? PI_MEM_TYPE_IMAGE3D + : (Desc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(Desc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(Desc.channel_order); + + MImpl->MSrcOffset = {0, 0, 0}; + MImpl->MDestOffset = {0, 0, 0}; + MImpl->MCopyExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MHostExtent = {Desc.width, Desc.height, Desc.depth}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageDesc.image_row_pitch = Pitch; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = detail::getPiImageCopyFlags( + get_pointer_type(Src, MQueue->get_context()), + get_pointer_type(Dest, MQueue->get_context())); + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_copy( + void *Src, sycl::range<3> SrcOffset, void *Dest, sycl::range<3> DestOffset, + const ext::oneapi::experimental::image_descriptor &DeviceImgDesc, + size_t DeviceRowPitch, sycl::range<3> HostExtent, + sycl::range<3> CopyExtent) { + MSrcPtr = Src; + MDstPtr = Dest; + + sycl::detail::pi::PiMemImageDesc PiDesc = {}; + PiDesc.image_width = DeviceImgDesc.width; + PiDesc.image_height = DeviceImgDesc.height; + PiDesc.image_depth = DeviceImgDesc.depth; + PiDesc.image_type = DeviceImgDesc.depth > 0 + ? PI_MEM_TYPE_IMAGE3D + : (DeviceImgDesc.height > 0 ? PI_MEM_TYPE_IMAGE2D + : PI_MEM_TYPE_IMAGE1D); + + sycl::detail::pi::PiMemImageFormat PiFormat; + PiFormat.image_channel_data_type = + sycl::_V1::detail::convertChannelType(DeviceImgDesc.channel_type); + PiFormat.image_channel_order = + sycl::_V1::detail::convertChannelOrder(DeviceImgDesc.channel_order); + + MImpl->MSrcOffset = {SrcOffset[0], SrcOffset[1], SrcOffset[2]}; + MImpl->MDestOffset = {DestOffset[0], DestOffset[1], DestOffset[2]}; + MImpl->MHostExtent = {HostExtent[0], HostExtent[1], HostExtent[2]}; + MImpl->MCopyExtent = {CopyExtent[0], CopyExtent[1], CopyExtent[2]}; + MImpl->MImageDesc = PiDesc; + MImpl->MImageDesc.image_row_pitch = DeviceRowPitch; + MImpl->MImageFormat = PiFormat; + MImpl->MImageCopyFlags = detail::getPiImageCopyFlags( + get_pointer_type(Src, MQueue->get_context()), + get_pointer_type(Dest, MQueue->get_context())); + setType(detail::CG::CopyImage); +} + +void handler::ext_oneapi_wait_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) { + MImpl->MInteropSemaphoreHandle = + (sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle; + setType(detail::CG::SemaphoreWait); +} + +void handler::ext_oneapi_signal_external_semaphore( + sycl::ext::oneapi::experimental::interop_semaphore_handle SemaphoreHandle) { + MImpl->MInteropSemaphoreHandle = + (sycl::detail::pi::PiInteropSemaphoreHandle)SemaphoreHandle.raw_handle; + setType(detail::CG::SemaphoreSignal); +} + void handler::use_kernel_bundle( const kernel_bundle &ExecBundle) { diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index f15400d55faf8..c5b1a56dae6d4 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3661,10 +3661,58 @@ _ZN4sycl3_V13ext6oneapi10level_zero10make_queueERKNS0_7contextERKNS0_6deviceEmbb _ZN4sycl3_V13ext6oneapi10level_zero11make_deviceERKNS0_8platformEm _ZN4sycl3_V13ext6oneapi10level_zero12make_contextERKSt6vectorINS0_6deviceESaIS5_EEmb _ZN4sycl3_V13ext6oneapi10level_zero13make_platformEm +_ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageENS3_16image_mem_handleERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageEPvmRKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageEPvmRKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental12create_imageERNS3_9image_memERKNS3_22bindless_image_samplerERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental14free_image_memENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental15alloc_image_memERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental15alloc_image_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental15free_mipmap_memENS3_16image_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental15free_mipmap_memENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental15get_image_rangeENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental16alloc_mipmap_memERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental16alloc_mipmap_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_20sampled_image_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_22unsampled_image_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental20destroy_image_handleERNS3_22unsampled_image_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmRKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental20pitched_alloc_deviceEPmmmjRKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental22get_image_channel_typeENS3_16image_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental22get_image_channel_typeENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental22get_image_num_channelsENS3_16image_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental22get_image_num_channelsENS3_16image_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental22import_external_memoryINS3_15external_mem_fdEEENS3_18interop_mem_handleENS3_23external_mem_descriptorIT_EERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental22import_external_memoryINS3_15external_mem_fdEEENS3_18interop_mem_handleENS3_23external_mem_descriptorIT_EERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental23prepare_for_device_copyEPKvmRKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental23prepare_for_device_copyEPKvmRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental23release_external_memoryENS3_18interop_mem_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental23release_external_memoryENS3_18interop_mem_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental24get_mip_level_mem_handleENS3_16image_mem_handleEjRKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental24get_mip_level_mem_handleENS3_16image_mem_handleEjRKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental24release_from_device_copyEPKvRKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental24release_from_device_copyEPKvRKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21external_semaphore_fdEEENS3_24interop_semaphore_handleENS3_29external_semaphore_descriptorIT_EERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21external_semaphore_fdEEENS3_24interop_semaphore_handleENS3_29external_semaphore_descriptorIT_EERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_mem_handleERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental25map_external_memory_arrayENS3_18interop_mem_handleERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental26destroy_external_semaphoreENS3_24interop_semaphore_handleERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implC2ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD1Ev +_ZN4sycl3_V13ext6oneapi12experimental6detail14image_mem_implD2Ev _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph12finalizeImplEv _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graph6updateERKNS3_13command_graphILNS3_11graph_stateE0EEE _ZN4sycl3_V13ext6oneapi12experimental6detail24executable_command_graphC1ERKSt10shared_ptrINS4_10graph_implEERKNS0_7contextE @@ -3679,6 +3727,10 @@ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph7addImplES _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph9make_edgeERNS3_4nodeES7_ _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC1ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE _ZN4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graphC2ERKNS0_7contextERKNS0_6deviceERKNS0_13property_listE +_ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental9image_memC1ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental9image_memC2ERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental9image_memC2ERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi15filter_selectorC1ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN4sycl3_V13ext6oneapi15filter_selectorC2ERKNSt7__cxx1112basic_stringIcSt11char_traitsIcESaIcEEE _ZN4sycl3_V13ext8codeplay12experimental14fusion_wrapper12start_fusionEv @@ -3850,6 +3902,7 @@ _ZN4sycl3_V16detail13MemoryManager16allocateMemImageESt10shared_ptrINS1_12contex _ZN4sycl3_V16detail13MemoryManager17allocateMemBufferESt10shared_ptrINS1_12context_implEEPNS1_11SYCLMemObjIEPvbmRKS3_INS1_10event_implEERKS5_RKNS0_13property_listERP9_pi_event _ZN4sycl3_V16detail13MemoryManager18allocateHostMemoryEPNS1_11SYCLMemObjIEPvbmRKNS0_13property_listE _ZN4sycl3_V16detail13MemoryManager19allocateImageObjectESt10shared_ptrINS1_12context_implEEPvbRK14_pi_image_descRK16_pi_image_formatRKNS0_13property_listE +_ZN4sycl3_V16detail13MemoryManager19copy_image_bindlessEPvSt10shared_ptrINS1_10queue_implEES3_RK14_pi_image_descRK16_pi_image_format20_pi_image_copy_flags22pi_image_offset_structSE_22pi_image_region_structSF_RKSt6vectorIP9_pi_eventSaISI_EEPSI_ _ZN4sycl3_V16detail13MemoryManager20allocateBufferObjectESt10shared_ptrINS1_12context_implEEPvbmRKNS0_13property_listE _ZN4sycl3_V16detail13MemoryManager20allocateMemSubBufferESt10shared_ptrINS1_12context_implEEPvmmNS0_5rangeILi3EEESt6vectorIS3_INS1_10event_implEESaISB_EERP9_pi_event _ZN4sycl3_V16detail13MemoryManager21copy_to_device_globalEPKvbSt10shared_ptrINS1_10queue_implEEmmS4_RKSt6vectorIP9_pi_eventSaISA_EEPSA_ @@ -4019,6 +4072,12 @@ _ZN4sycl3_V17handler10mem_adviseEPKvmi _ZN4sycl3_V17handler10processArgEPvRKNS0_6detail19kernel_param_kind_tEimRmbb _ZN4sycl3_V17handler12addReductionERKSt10shared_ptrIKvE _ZN4sycl3_V17handler13getKernelNameB5cxx11Ev +_ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleENS0_5rangeILi3EEERKNS4_16image_descriptorEPvS7_S7_S7_ +_ZN4sycl3_V17handler15ext_oneapi_copyENS0_3ext6oneapi12experimental16image_mem_handleEPvRKNS4_16image_descriptorE +_ZN4sycl3_V17handler15ext_oneapi_copyEPvNS0_3ext6oneapi12experimental16image_mem_handleERKNS5_16image_descriptorE +_ZN4sycl3_V17handler15ext_oneapi_copyEPvNS0_5rangeILi3EEES2_S4_RKNS0_3ext6oneapi12experimental16image_descriptorEmS4_S4_ +_ZN4sycl3_V17handler15ext_oneapi_copyEPvNS0_5rangeILi3EEES4_NS0_3ext6oneapi12experimental16image_mem_handleES4_RKNS7_16image_descriptorES4_ +_ZN4sycl3_V17handler15ext_oneapi_copyEPvS2_RKNS0_3ext6oneapi12experimental16image_descriptorEm _ZN4sycl3_V17handler16ext_oneapi_graphENS0_3ext6oneapi12experimental13command_graphILNS4_11graph_stateE1EEE _ZN4sycl3_V17handler17supportsUSMFill2DEv _ZN4sycl3_V17handler17use_kernel_bundleERKNS0_13kernel_bundleILNS0_12bundle_stateE2EEE @@ -4050,6 +4109,8 @@ _ZN4sycl3_V17handler28extractArgsAndReqsFromLambdaEPcmPKNS0_6detail19kernel_para _ZN4sycl3_V17handler28memcpyToHostOnlyDeviceGlobalEPKvS3_mbmm _ZN4sycl3_V17handler28setStateExplicitKernelBundleEv _ZN4sycl3_V17handler30memcpyFromHostOnlyDeviceGlobalEPvPKvbmm +_ZN4sycl3_V17handler34ext_oneapi_wait_external_semaphoreENS0_3ext6oneapi12experimental24interop_semaphore_handleE +_ZN4sycl3_V17handler36ext_oneapi_signal_external_semaphoreENS0_3ext6oneapi12experimental24interop_semaphore_handleE _ZN4sycl3_V17handler6memcpyEPvPKvm _ZN4sycl3_V17handler6memsetEPvim _ZN4sycl3_V17handler8finalizeEv @@ -4119,6 +4180,12 @@ _ZNK4sycl3_V115device_selector13select_deviceEv _ZNK4sycl3_V116default_selectorclERKNS0_6deviceE _ZNK4sycl3_V120accelerator_selectorclERKNS0_6deviceE _ZNK4sycl3_V13ext6oneapi12experimental6detail24modifiable_command_graph8finalizeERKNS0_13property_listE +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem16get_channel_typeEv +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem16get_num_channelsEv +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem17get_channel_orderEv +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem24get_mip_level_mem_handleEj +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem8get_typeEv +_ZNK4sycl3_V13ext6oneapi12experimental9image_mem9get_rangeEv _ZNK4sycl3_V13ext6oneapi15filter_selector13select_deviceEv _ZNK4sycl3_V13ext6oneapi15filter_selector5resetEv _ZNK4sycl3_V13ext6oneapi15filter_selectorclERKNS0_6deviceE @@ -4198,7 +4265,12 @@ _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6de _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi3EEEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device21image_row_pitch_alignEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device21mipmap_max_anisotropyEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device22max_global_work_groupsEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device22max_image_linear_widthEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device23max_image_linear_heightEEENT_11return_typeEv +_ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext6oneapi12experimental4info6device26max_image_linear_row_pitchEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext8codeplay12experimental4info6device15supports_fusionEEENT_11return_typeEv _ZNK4sycl3_V16detail11device_impl8get_infoINS0_3ext8codeplay12experimental4info6device28max_registers_per_work_groupEEENT_11return_typeEv _ZNK4sycl3_V16detail11image_plain10getSamplerEv @@ -4371,7 +4443,12 @@ _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device13graph_s _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi1EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi2EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device15max_work_groupsILi3EEEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device21image_row_pitch_alignEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device21mipmap_max_anisotropyEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device22max_global_work_groupsEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device22max_image_linear_widthEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device23max_image_linear_heightEEENS0_6detail19is_device_info_descIT_E11return_typeEv +_ZNK4sycl3_V16device8get_infoINS0_3ext6oneapi12experimental4info6device26max_image_linear_row_pitchEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext8codeplay12experimental4info6device15supports_fusionEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_3ext8codeplay12experimental4info6device28max_registers_per_work_groupEEENS0_6detail19is_device_info_descIT_E11return_typeEv _ZNK4sycl3_V16device8get_infoINS0_4info6device10extensionsEEENS0_6detail19is_device_info_descIT_E11return_typeEv diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 28923141bfe63..b9182a5f7c8c7 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -110,6 +110,8 @@ ??$get_info@Uimage3d_max_width@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Uimage_max_array_size@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Uimage_max_buffer_size@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ +??$get_info@Uimage_row_pitch_align@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ +??$get_info@Uimage_row_pitch_align@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Uimage_support@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ ??$get_info@Uis_available@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ ??$get_info@Uis_compiler_available@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_NXZ @@ -126,6 +128,12 @@ ??$get_info@Umax_constant_buffer_size@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Umax_global_work_groups@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Umax_global_work_groups@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA_KXZ +??$get_info@Umax_image_linear_height@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ +??$get_info@Umax_image_linear_height@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ +??$get_info@Umax_image_linear_row_pitch@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ +??$get_info@Umax_image_linear_row_pitch@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ +??$get_info@Umax_image_linear_width@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ +??$get_info@Umax_image_linear_width@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Umax_mem_alloc_size@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Umax_mem_bandwidth@device@info@intel@ext@_V1@sycl@@@device@_V1@sycl@@QEBA_KXZ ??$get_info@Umax_mem_bandwidth@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBA_KXZ @@ -146,6 +154,8 @@ ??$get_info@Umemory_bus_width@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ ??$get_info@Umemory_clock_rate@device@info@intel@ext@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ ??$get_info@Umemory_clock_rate@device@info@intel@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAIXZ +??$get_info@Umipmap_max_anisotropy@device@info@experimental@oneapi@ext@_V1@sycl@@@device@_V1@sycl@@QEBAMXZ +??$get_info@Umipmap_max_anisotropy@device@info@experimental@oneapi@ext@_V1@sycl@@@device_impl@detail@_V1@sycl@@QEBAMXZ ??$get_info@Uname@device@info@_V1@sycl@@@device@_V1@sycl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ ??$get_info@Uname@platform@info@_V1@sycl@@@platform@_V1@sycl@@QEBA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@XZ ??$get_info@Unative_vector_width_char@device@info@_V1@sycl@@@device@_V1@sycl@@QEBAIXZ @@ -400,6 +410,10 @@ ??$has_property@Vuse_primary_context@cuda@context@property@_V1@sycl@@@image_plain@detail@_V1@sycl@@IEBA_NXZ ??$has_property@Vuse_primary_context@cuda@context@property@_V1@sycl@@@sampler@_V1@sycl@@QEBA_NXZ ??$has_property@Vuse_primary_context@cuda@context@property@_V1@sycl@@@stream@_V1@sycl@@QEBA_NXZ +??$import_external_memory@Uexternal_mem_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUinterop_mem_handle@01234@U?$external_mem_descriptor@Uexternal_mem_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z +??$import_external_memory@Uexternal_mem_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUinterop_mem_handle@01234@U?$external_mem_descriptor@Uexternal_mem_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z +??$import_external_semaphore@Uexternal_semaphore_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUinterop_semaphore_handle@01234@U?$external_semaphore_descriptor@Uexternal_semaphore_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z +??$import_external_semaphore@Uexternal_semaphore_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUinterop_semaphore_handle@01234@U?$external_semaphore_descriptor@Uexternal_semaphore_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??0AccessorBaseHost@detail@_V1@sycl@@IEAA@AEBV?$shared_ptr@VAccessorImplHost@detail@_V1@sycl@@@std@@@Z ??0AccessorBaseHost@detail@_V1@sycl@@QEAA@$$QEAV0123@@Z ??0AccessorBaseHost@detail@_V1@sycl@@QEAA@AEBV0123@@Z @@ -543,6 +557,12 @@ ??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_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@$$QEAV012345@@Z +??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBV012345@@Z +??0image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ +??0image_mem_impl@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@AEBUimage_descriptor@23456@AEBVdevice@56@AEBVcontext@56@@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@AEBV?$shared_ptr@$$CBX@std@@W4image_channel_order@23@W4image_channel_type@23@Uimage_sampler@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@@Z @@ -643,6 +663,8 @@ ??1handler@_V1@sycl@@AEAA@XZ ??1host_selector@_V1@sycl@@UEAA@XZ ??1image_impl@detail@_V1@sycl@@UEAA@XZ +??1image_mem@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ +??1image_mem_impl@detail@experimental@oneapi@ext@_V1@sycl@@QEAA@XZ ??1image_plain@detail@_V1@sycl@@QEAA@XZ ??1kernel@_V1@sycl@@QEAA@XZ ??1kernel_bundle_plain@detail@_V1@sycl@@QEAA@XZ @@ -722,6 +744,8 @@ ??4half@host_half_impl@detail@_V1@sycl@@QEAAAEAV01234@AEBV01234@@Z ??4host_selector@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z ??4host_selector@_V1@sycl@@QEAAAEAV012@AEBV012@@Z +??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@$$QEAV012345@@Z +??4image_mem@experimental@oneapi@ext@_V1@sycl@@QEAAAEAV012345@AEBV012345@@Z ??4image_plain@detail@_V1@sycl@@QEAAAEAV0123@$$QEAV0123@@Z ??4image_plain@detail@_V1@sycl@@QEAAAEAV0123@AEBV0123@@Z ??4kernel@_V1@sycl@@QEAAAEAV012@$$QEAV012@@Z @@ -749,6 +773,7 @@ ??8device@_V1@sycl@@QEBA_NAEBV012@@Z ??8device_image_plain@detail@_V1@sycl@@QEBA_NAEBV0123@@Z ??8event@_V1@sycl@@QEBA_NAEBV012@@Z +??8image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA_NAEBV012345@@Z ??8kernel@_V1@sycl@@QEBA_NAEBV012@@Z ??8kernel_bundle_plain@detail@_V1@sycl@@QEBA_NAEBV0123@@Z ??8kernel_id@_V1@sycl@@QEBA_NAEBV012@@Z @@ -760,6 +785,7 @@ ??9device@_V1@sycl@@QEBA_NAEBV012@@Z ??9device_image_plain@detail@_V1@sycl@@QEBA_NAEBV0123@@Z ??9event@_V1@sycl@@QEBA_NAEBV012@@Z +??9image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA_NAEBV012345@@Z ??9kernel@_V1@sycl@@QEBA_NAEBV012@@Z ??9kernel_bundle_plain@detail@_V1@sycl@@QEBA_NAEBV0123@@Z ??9kernel_id@_V1@sycl@@QEBA_NAEBV012@@Z @@ -870,6 +896,10 @@ ?aligned_alloc_shared@_V1@sycl@@YAPEAX_K0AEBVdevice@12@AEBVcontext@12@AEBVproperty_list@12@AEBUcode_location@detail@12@@Z ?aligned_alloc_shared@_V1@sycl@@YAPEAX_K0AEBVqueue@12@AEBUcode_location@detail@12@@Z ?aligned_alloc_shared@_V1@sycl@@YAPEAX_K0AEBVqueue@12@AEBVproperty_list@12@AEBUcode_location@detail@12@@Z +?alloc_image_mem@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?alloc_image_mem@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?alloc_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?alloc_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z ?allocate@MemoryManager@detail@_V1@sycl@@SAPEAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAVSYCLMemObjI@234@_NPEAXV?$vector@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@V?$allocator@V?$shared_ptr@Vevent_impl@detail@_V1@sycl@@@std@@@2@@6@AEAPEAU_pi_event@@@Z ?allocateBufferObject@MemoryManager@detail@_V1@sycl@@SAPEAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAX_N_KAEBVproperty_list@34@@Z ?allocateHostMem@SYCLMemObjT@detail@_V1@sycl@@UEAAPEAXXZ @@ -917,14 +947,31 @@ ?copy@MemoryManager@detail@_V1@sycl@@SAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@IV?$range@$02@34@3V?$id@$02@34@I12I334IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z ?copy_2d_usm@MemoryManager@detail@_V1@sycl@@SAXPEBX_KV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@PEAX111V?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z ?copy_from_device_global@MemoryManager@detail@_V1@sycl@@SAXPEBX_NV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_K3PEAXAEBV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z +?copy_image_bindless@MemoryManager@detail@_V1@sycl@@SAXPEAXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@0AEBU_pi_image_desc@@AEBU_pi_image_format@@W4_pi_image_copy_flags@@Upi_image_offset_struct@@5Upi_image_region_struct@@6AEBV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z ?copy_to_device_global@MemoryManager@detail@_V1@sycl@@SAXPEBX_NV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_K30AEBV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z ?copy_usm@MemoryManager@detail@_V1@sycl@@SAXPEBXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_KPEAXV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@6@PEAPEAU_pi_event@@@Z ?cpu_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@AEAVimage_mem@12345@AEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@AEAVimage_mem@12345@AEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@PEAX_KAEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@PEAX_KAEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUbindless_image_sampler@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@AEAVimage_mem@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@AEAVimage_mem@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?create_image@experimental@oneapi@ext@_V1@sycl@@YA?AUunsampled_image_handle@12345@Uimage_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z ?default_selector_v@_V1@sycl@@YAHAEBVdevice@12@@Z ?deleteAccProps@buffer_plain@detail@_V1@sycl@@IEAAXAEBW4PropWithDataKind@234@@Z ?deleteAccessorProperty@SYCLMemObjT@detail@_V1@sycl@@QEAAXAEBW4PropWithDataKind@234@@Z ?depends_on@handler@_V1@sycl@@QEAAXAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@@Z ?depends_on@handler@_V1@sycl@@QEAAXVevent@23@@Z +?destroy_external_semaphore@experimental@oneapi@ext@_V1@sycl@@YAXUinterop_semaphore_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?destroy_external_semaphore@experimental@oneapi@ext@_V1@sycl@@YAXUinterop_semaphore_handle@12345@AEBVqueue@45@@Z +?destroy_image_handle@experimental@oneapi@ext@_V1@sycl@@YAXAEAUsampled_image_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?destroy_image_handle@experimental@oneapi@ext@_V1@sycl@@YAXAEAUsampled_image_handle@12345@AEBVqueue@45@@Z +?destroy_image_handle@experimental@oneapi@ext@_V1@sycl@@YAXAEAUunsampled_image_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?destroy_image_handle@experimental@oneapi@ext@_V1@sycl@@YAXAEAUunsampled_image_handle@12345@AEBVqueue@45@@Z ?destructorNotification@buffer_impl@detail@_V1@sycl@@QEAAXPEAX@Z ?detachMemoryObject@SYCLMemObjT@detail@_V1@sycl@@QEBAXAEBV?$shared_ptr@VSYCLMemObjT@detail@_V1@sycl@@@std@@@Z ?determineHostPtr@SYCLMemObjT@detail@_V1@sycl@@IEAAXAEBV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@_NAEAPEAXAEA_N@Z @@ -945,6 +992,30 @@ ?ext_oneapi_barrier@handler@_V1@sycl@@QEAAXAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@@Z ?ext_oneapi_barrier@handler@_V1@sycl@@QEAAXXZ ?ext_oneapi_can_access_peer@device@_V1@sycl@@QEAA_NAEBV123@W4peer_access@oneapi@ext@23@@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXPEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_K@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXPEAXUimage_mem_handle@experimental@oneapi@ext@23@AEBUimage_descriptor@56723@@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXPEAXV?$range@$02@23@01AEBUimage_descriptor@experimental@oneapi@ext@23@_K11@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXPEAXV?$range@$02@23@1Uimage_mem_handle@experimental@oneapi@ext@23@1AEBUimage_descriptor@67823@1@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@56723@@Z +?ext_oneapi_copy@handler@_V1@sycl@@QEAAXUimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@56723@PEAX111@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAX0AEBUimage_descriptor@experimental@oneapi@ext@23@_KV423@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXUimage_mem_handle@experimental@oneapi@ext@23@AEBUimage_descriptor@67823@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXUimage_mem_handle@experimental@oneapi@ext@23@AEBUimage_descriptor@67823@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXUimage_mem_handle@experimental@oneapi@ext@23@AEBUimage_descriptor@67823@V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@01AEBUimage_descriptor@experimental@oneapi@ext@23@_K11AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@01AEBUimage_descriptor@experimental@oneapi@ext@23@_K11AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@01AEBUimage_descriptor@experimental@oneapi@ext@23@_K11V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@1Uimage_mem_handle@experimental@oneapi@ext@23@1AEBUimage_descriptor@78923@1AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@1Uimage_mem_handle@experimental@oneapi@ext@23@1AEBUimage_descriptor@78923@1AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@PEAXV?$range@$02@23@1Uimage_mem_handle@experimental@oneapi@ext@23@1AEBUimage_descriptor@78923@1V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@67823@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@67823@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@PEAXAEBUimage_descriptor@67823@V423@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@67823@PEAX111AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@67823@PEAX111AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_copy@queue@_V1@sycl@@QEAA?AVevent@23@Uimage_mem_handle@experimental@oneapi@ext@23@V?$range@$02@23@AEBUimage_descriptor@67823@PEAX111V423@AEBUcode_location@detail@23@@Z ?ext_oneapi_copyD2D_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAVSYCLMemObjI@234@PEAXIV?$range@$02@34@4V?$id@$02@34@I3I445IV?$vector@IV?$allocator@I@std@@@6@PEAI@Z ?ext_oneapi_copyD2H_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAVSYCLMemObjI@234@PEAXIV?$range@$02@34@4V?$id@$02@34@IPEADI45IV?$vector@IV?$allocator@I@std@@@6@PEAI@Z ?ext_oneapi_copyH2D_cmd_buffer@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAU_pi_ext_command_buffer@@PEAVSYCLMemObjI@234@PEADIV?$range@$02@34@V?$id@$02@34@IPEAXI445IV?$vector@IV?$allocator@I@std@@@6@PEAI@Z @@ -977,8 +1048,16 @@ ?ext_oneapi_owner_before@?$OwnerLessBase@Vqueue@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVqueue@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBV?$weak_object_base@Vstream@_V1@sycl@@@2oneapi@ext@34@@Z ?ext_oneapi_owner_before@?$OwnerLessBase@Vstream@_V1@sycl@@@detail@_V1@sycl@@QEBA_NAEBVstream@34@@Z +?ext_oneapi_signal_external_semaphore@handler@_V1@sycl@@QEAAXUinterop_semaphore_handle@experimental@oneapi@ext@23@@Z +?ext_oneapi_signal_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z +?ext_oneapi_signal_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_signal_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@V423@AEBUcode_location@detail@23@@Z ?ext_oneapi_submit_barrier@queue@_V1@sycl@@QEAA?AVevent@23@AEBUcode_location@detail@23@@Z ?ext_oneapi_submit_barrier@queue@_V1@sycl@@QEAA?AVevent@23@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_wait_external_semaphore@handler@_V1@sycl@@QEAAXUinterop_semaphore_handle@experimental@oneapi@ext@23@@Z +?ext_oneapi_wait_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@AEBUcode_location@detail@23@@Z +?ext_oneapi_wait_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@AEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z +?ext_oneapi_wait_external_semaphore@queue@_V1@sycl@@QEAA?AVevent@23@Uinterop_semaphore_handle@experimental@oneapi@ext@23@V423@AEBUcode_location@detail@23@@Z ?extractArgsAndReqs@handler@_V1@sycl@@AEAAXXZ ?extractArgsAndReqsFromLambda@handler@_V1@sycl@@AEAAXPEAD_KPEBUkernel_param_desc_t@detail@23@_N@Z ?fill@MemoryManager@detail@_V1@sycl@@SAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@_KPEBDIV?$range@$02@34@5V?$id@$02@34@IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z @@ -992,6 +1071,10 @@ ?flush@stream_impl@detail@_V1@sycl@@QEAAXXZ ?free@_V1@sycl@@YAXPEAXAEBVcontext@12@AEBUcode_location@detail@12@@Z ?free@_V1@sycl@@YAXPEAXAEBVqueue@12@AEBUcode_location@detail@12@@Z +?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?free_image_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z +?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?free_mipmap_mem@experimental@oneapi@ext@_V1@sycl@@YAXUimage_mem_handle@12345@AEBVqueue@45@@Z ?get@context@_V1@sycl@@QEBAPEAU_cl_context@@XZ ?get@device@_V1@sycl@@QEBAPEAU_cl_device_id@@XZ ?get@kernel@_V1@sycl@@QEBAPEAU_cl_kernel@@XZ @@ -1112,8 +1195,11 @@ ?get_backend@kernel_bundle_plain@detail@_V1@sycl@@QEBA?AW4backend@34@XZ ?get_backend@platform@_V1@sycl@@QEBA?AW4backend@23@XZ ?get_backend@queue@_V1@sycl@@QEBA?AW4backend@23@XZ +?get_channel_order@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_channel_order@56@XZ +?get_channel_type@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_channel_type@56@XZ ?get_cl_code@exception@_V1@sycl@@QEBAHXZ ?get_context@exception@_V1@sycl@@QEBA?AVcontext@23@XZ +?get_context@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVcontext@56@XZ ?get_context@kernel@_V1@sycl@@QEBA?AVcontext@23@XZ ?get_context@kernel_bundle_plain@detail@_V1@sycl@@QEBA?AVcontext@34@XZ ?get_context@queue@_V1@sycl@@QEBA?AVcontext@23@XZ @@ -1122,6 +1208,8 @@ ?get_count@SYCLMemObjT@detail@_V1@sycl@@QEBA_KXZ ?get_count@image_impl@detail@_V1@sycl@@QEBA_KXZ ?get_count@image_plain@detail@_V1@sycl@@IEBA_KXZ +?get_descriptor@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBAAEBUimage_descriptor@23456@XZ +?get_device@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AVdevice@56@XZ ?get_device@queue@_V1@sycl@@QEBA?AVdevice@23@XZ ?get_devices@context@_V1@sycl@@QEBA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@XZ ?get_devices@device@_V1@sycl@@SA?AV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@W4device_type@info@23@@Z @@ -1131,6 +1219,13 @@ ?get_filtering_mode@sampler@_V1@sycl@@QEBA?AW4filtering_mode@23@XZ ?get_filtering_mode@sampler_impl@detail@_V1@sycl@@QEBA?AW4filtering_mode@34@XZ ?get_flags@stream@_V1@sycl@@AEBAIXZ +?get_handle@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AUimage_mem_handle@23456@XZ +?get_image_channel_type@experimental@oneapi@ext@_V1@sycl@@YA?AW4image_channel_type@45@Uimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?get_image_channel_type@experimental@oneapi@ext@_V1@sycl@@YA?AW4image_channel_type@45@Uimage_mem_handle@12345@AEBVqueue@45@@Z +?get_image_num_channels@experimental@oneapi@ext@_V1@sycl@@YAIUimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?get_image_num_channels@experimental@oneapi@ext@_V1@sycl@@YAIUimage_mem_handle@12345@AEBVqueue@45@@Z +?get_image_range@experimental@oneapi@ext@_V1@sycl@@YA?AV?$range@$02@45@Uimage_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?get_image_range@experimental@oneapi@ext@_V1@sycl@@YA?AV?$range@$02@45@Uimage_mem_handle@12345@AEBVqueue@45@@Z ?get_kernel@kernel_bundle_plain@detail@_V1@sycl@@IEBA?AVkernel@34@AEBVkernel_id@34@@Z ?get_kernel_bundle@kernel@_V1@sycl@@QEBA?AV?$kernel_bundle@$01@23@XZ ?get_kernel_bundle_impl@detail@_V1@sycl@@YA?AV?$shared_ptr@Vkernel_bundle_impl@detail@_V1@sycl@@@std@@AEBVcontext@23@AEBV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@5@AEBV?$vector@Vkernel_id@_V1@sycl@@V?$allocator@Vkernel_id@_V1@sycl@@@std@@@5@W4bundle_state@23@@Z @@ -1141,7 +1236,11 @@ ?get_kernel_ids@kernel_bundle_plain@detail@_V1@sycl@@QEBA?AV?$vector@Vkernel_id@_V1@sycl@@V?$allocator@Vkernel_id@_V1@sycl@@@std@@@std@@XZ ?get_max_statement_size@stream@_V1@sycl@@QEBA_KXZ ?get_max_statement_size@stream_impl@detail@_V1@sycl@@QEBA_KXZ +?get_mip_level_mem_handle@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@U612345@IAEBVdevice@45@AEBVcontext@45@@Z +?get_mip_level_mem_handle@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@U612345@IAEBVqueue@45@@Z +?get_mip_level_mem_handle@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AUimage_mem_handle@23456@I@Z ?get_name@kernel_id@_V1@sycl@@QEBAPEBDXZ +?get_num_channels@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBAIXZ ?get_pipe_name@pipe_base@experimental@intel@ext@_V1@sycl@@KA?AV?$basic_string@DU?$char_traits@D@std@@V?$allocator@D@2@@std@@PEBX@Z ?get_pitch@image_impl@detail@_V1@sycl@@QEBA?AV?$range@$01@34@XZ ?get_pitch@image_plain@detail@_V1@sycl@@IEBA?AV?$range@$01@34@XZ @@ -1153,12 +1252,14 @@ ?get_precision@stream@_V1@sycl@@QEBA_KXZ ?get_queue@fusion_wrapper@experimental@codeplay@ext@_V1@sycl@@QEBA?AVqueue@56@XZ ?get_range@image_impl@detail@_V1@sycl@@QEBA?AV?$range@$02@34@XZ +?get_range@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AV?$range@$02@56@XZ ?get_range@image_plain@detail@_V1@sycl@@IEBA?AV?$range@$02@34@XZ ?get_size@image_plain@detail@_V1@sycl@@IEBA_KXZ ?get_size@stream@_V1@sycl@@QEBA_KXZ ?get_size@stream_impl@detail@_V1@sycl@@QEBA_KXZ ?get_specialization_constant_impl@kernel_bundle_plain@detail@_V1@sycl@@IEBAXPEBDPEAX@Z ?get_stream_mode@stream@_V1@sycl@@QEBA?AW4stream_manipulator@23@XZ +?get_type@image_mem@experimental@oneapi@ext@_V1@sycl@@QEBA?AW4image_type@23456@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@_V1@sycl@@QEBA_KXZ @@ -1253,6 +1354,8 @@ ?malloc_shared@_V1@sycl@@YAPEAX_KAEBVqueue@12@AEBUcode_location@detail@12@@Z ?malloc_shared@_V1@sycl@@YAPEAX_KAEBVqueue@12@AEBVproperty_list@12@AEBUcode_location@detail@12@@Z ?map@MemoryManager@detail@_V1@sycl@@SAPEAXPEAVSYCLMemObjI@234@PEAXV?$shared_ptr@Vqueue_impl@detail@_V1@sycl@@@std@@W4mode@access@34@IV?$range@$02@34@4V?$id@$02@34@IV?$vector@PEAU_pi_event@@V?$allocator@PEAU_pi_event@@@std@@@7@AEAPEAU_pi_event@@@Z +?map_external_memory_array@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@Uinterop_mem_handle@12345@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?map_external_memory_array@experimental@oneapi@ext@_V1@sycl@@YA?AUimage_mem_handle@12345@Uinterop_mem_handle@12345@AEBUimage_descriptor@12345@AEBVqueue@45@@Z ?markAsInternal@SYCLMemObjT@detail@_V1@sycl@@QEAAXXZ ?markBufferAsInternal@detail@_V1@sycl@@YAXAEBV?$shared_ptr@Vbuffer_impl@detail@_V1@sycl@@@std@@@Z ?mem_advise@handler@_V1@sycl@@QEAAXPEBX_KH@Z @@ -1281,6 +1384,10 @@ ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$00@23@Vkernel@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$01@23@Vkernel@23@@Z ?parallel_for@handler@_V1@sycl@@QEAAXV?$range@$02@23@Vkernel@23@@Z +?pitched_alloc_device@experimental@oneapi@ext@_V1@sycl@@YAPEAXPEA_KAEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?pitched_alloc_device@experimental@oneapi@ext@_V1@sycl@@YAPEAXPEA_KAEBUimage_descriptor@12345@AEBVqueue@45@@Z +?pitched_alloc_device@experimental@oneapi@ext@_V1@sycl@@YAPEAXPEA_K_K1IAEBVdevice@45@AEBVcontext@45@@Z +?pitched_alloc_device@experimental@oneapi@ext@_V1@sycl@@YAPEAXPEA_K_K1IAEBVqueue@45@@Z ?prefetch@handler@_V1@sycl@@QEAAXPEBX_K@Z ?prefetch@queue@_V1@sycl@@QEAA?AVevent@23@PEBX_KAEBUcode_location@detail@23@@Z ?prefetch@queue@_V1@sycl@@QEAA?AVevent@23@PEBX_KAEBV?$vector@Vevent@_V1@sycl@@V?$allocator@Vevent@_V1@sycl@@@std@@@std@@AEBUcode_location@detail@23@@Z @@ -1298,6 +1405,8 @@ ?releaseHostMem@SYCLMemObjT@detail@_V1@sycl@@UEAAXPEAX@Z ?releaseMem@SYCLMemObjT@detail@_V1@sycl@@UEAAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAX@Z ?releaseMemObj@MemoryManager@detail@_V1@sycl@@SAXV?$shared_ptr@Vcontext_impl@detail@_V1@sycl@@@std@@PEAVSYCLMemObjI@234@PEAX2@Z +?release_external_memory@experimental@oneapi@ext@_V1@sycl@@YAXUinterop_mem_handle@12345@AEBVdevice@45@AEBVcontext@45@@Z +?release_external_memory@experimental@oneapi@ext@_V1@sycl@@YAXUinterop_mem_handle@12345@AEBVqueue@45@@Z ?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVcontext@45@@Z ?release_from_device_copy@experimental@oneapi@ext@_V1@sycl@@YAXPEBXAEBVqueue@45@@Z ?removeDuplicateDevices@detail@_V1@sycl@@YA?BV?$vector@Vdevice@_V1@sycl@@V?$allocator@Vdevice@_V1@sycl@@@std@@@std@@AEBV45@@Z