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