diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index 5c0c547edb1cf..5c63e54a5b744 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -259,6 +259,78 @@ type of the returned `image_descriptor` will be `image_type::standard`. Only array image types support more than one array layer. +==== Querying image support + +Not all devices support all combinations of image channel type, the number of +channels, the type of backing memory, and dimensionality. We provide functions +to query device support for the allocation and creation of images for a given +`image_descriptor` and the type of backing memory. + +===== Querying image memory support + +Before allocating memory for an image, the user may first query whether their +desired image backing memory type is supported by the device. + +The following query returns a vector of supported `image_memory_handle_type`s +based on the properties of a given `image_descriptor`. + +The `image_memory_handle_type::usm_pointer` relates to USM allocations, while +the `image_memory_handle_type::opaque_handle` relates to memory allocations of +the `image_mem_handle` type. + +If the returned vector is empty, this indicates that the device does not support +allocating or creating images for the specified `image_descriptor`. + +```cpp +namespace sycl::ext::oneapi::experimental { + +enum class image_memory_handle_type : /* unspecified */ { + usm_pointer, + opaque_handle +}; + +std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::queue &syclQueue); +} +``` + +===== Querying image handle support + +In order to query what types of image handles are supported for a combination +of a given `image_descriptor` and `image_memory_handle_type`, the user should +use the `is_image_handle_supported` query. + +The template parameter passed to this query should be either +`unsampled_image_handle` or `sampled_image_handle`. + +The boolean value returned from the query indicates whether the device supports +creating the given image handle type (sampled or unsampled) given the specified +`image_descriptor` and `image_memory_handle_type`. + +```cpp +namespace sycl::ext::oneapi::experimental { + +template +bool +is_image_handle_supported(const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +template +bool +is_image_handle_supported(const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::queue &syclQueue); +} +``` + === Allocating image memory The process of creating an image is two-fold: diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index b83e01b0bba9f..f49a4c1473260 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -593,6 +593,89 @@ __SYCL_EXPORT unsigned int get_image_num_channels(const image_mem_handle memHandle, const sycl::queue &syclQueue); +/** + * @brief Returns a vector of image-backing memory types supported by the + * device for a given `image_descriptor`. If the returned vector is + * empty, it indicates that the device does not support allocating or + * creating images with the properties described in the + * `image_descriptor`. + * + * @param imageDescriptor Properties of the image we want to query support + * for. + * @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 List of supported image-backing memory types + */ +__SYCL_EXPORT std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Returns a vector of image-backing memory types supported by the + * device for a given `image_descriptor`. If the returned vector is + * empty, it indicates that the device does not support allocating or + * creating images with the properties described in the + * `image_descriptor`. + * + * @param imageDescriptor Properties of the image we want to query support + * for. + * @param syclQueue The device/context association for which we want to query + * image memory support. + * @return List of supported image-backing memory types + */ +__SYCL_EXPORT std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::queue &syclQueue); + +/** + * @brief Returns `true` if the device supports creation of images of the + * ImageHandleType, given the combination of `image_descriptor` and + * `image_memory_handle_type`. + * + * @tparam ImageHandleType Either `sampled_image_handle` or + * `unsampled_image_handle`. + * @param imageDescriptor Properties of the image we want to query support + * for. + * @param imageMemoryHandleType Image memory handle type we want to query + * support for. + * @param syclDevice The device in which we want to query image handle + * support + * @param syclContext The context in which we want to query image handle + * support + * @return Boolean indicating support for image creation with the specified + * parameter. + */ + +template +__SYCL_EXPORT bool +is_image_handle_supported(const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::device &syclDevice, + const sycl::context &syclContext); + +/** + * @brief Returns `true` if the device supports creation of images of the + * ImageHandleType, given the combination of `image_descriptor` and + * `image_memory_handle_type`. + * + * @tparam ImageHandleType Either `sampled_image_handle` or + * `unsampled_image_handle` + * @param imageDescriptor Properties of the image we want to query support + * for. + * @param imageMemoryHandleType Image memory handle type we want to query + * support for. + * @param syclQueue The device/context association for which we want to query + * image handle support. + * @return Boolean indicating support for image creation with the specified + * parameter. + */ +template +__SYCL_EXPORT bool +is_image_handle_supported(const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::queue &syclQueue); + namespace detail { // is sycl::vec diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp index da3e254036baa..74e16d0c27be3 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images_memory.hpp @@ -107,6 +107,12 @@ enum image_copy_flags : unsigned int { DtoD = 2, }; +// The types of handles to image-backing memory +enum class image_memory_handle_type : unsigned int { + usm_pointer = 0, + opaque_handle = 1, +}; + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp index 95d68d49c3f16..fd4d812854851 100644 --- a/sycl/source/detail/bindless_images.cpp +++ b/sycl/source/detail/bindless_images.cpp @@ -793,6 +793,136 @@ get_image_num_channels(const image_mem_handle memHandle, syclQueue.get_context()); } +__SYCL_EXPORT std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::device &syclDevice, + const sycl::context &syclContext) { + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::AdapterPtr &Adapter = CtxImpl->getAdapter(); + + ur_image_desc_t urDesc; + ur_image_format_t urFormat; + populate_ur_structs(imageDescriptor, urDesc, urFormat); + + ur_bool_t supportsPointerAllocation{0}; + Adapter->call( + CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat, + ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_USM_POINTER, + &supportsPointerAllocation); + + ur_bool_t supportsOpaqueAllocation{0}; + Adapter->call( + CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat, + ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE, + &supportsOpaqueAllocation); + + std::vector supportedMemHandleTypes; + + if (supportsPointerAllocation) { + supportedMemHandleTypes.push_back(image_memory_handle_type::usm_pointer); + } + + if (supportsOpaqueAllocation) { + supportedMemHandleTypes.push_back(image_memory_handle_type::opaque_handle); + } + + return supportedMemHandleTypes; +} + +__SYCL_EXPORT std::vector +get_image_memory_support(const image_descriptor &imageDescriptor, + const sycl::queue &syclQueue) { + return get_image_memory_support(imageDescriptor, syclQueue.get_device(), + syclQueue.get_context()); +} + +template <> +__SYCL_EXPORT bool is_image_handle_supported( + const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::AdapterPtr &Adapter = CtxImpl->getAdapter(); + + ur_image_desc_t urDesc; + ur_image_format_t urFormat; + populate_ur_structs(imageDescriptor, urDesc, urFormat); + + const ur_exp_image_mem_type_t memHandleType = + (imageMemoryHandleType == image_memory_handle_type::opaque_handle) + ? ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE + : ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_USM_POINTER; + + ur_bool_t supportsUnsampledHandle{0}; + Adapter->call( + CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat, + memHandleType, &supportsUnsampledHandle); + + return supportsUnsampledHandle; +} + +template <> +__SYCL_EXPORT bool is_image_handle_supported( + const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::queue &syclQueue) { + return is_image_handle_supported( + imageDescriptor, imageMemoryHandleType, syclQueue.get_device(), + syclQueue.get_context()); +} + +template <> +__SYCL_EXPORT bool is_image_handle_supported( + const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::device &syclDevice, const sycl::context &syclContext) { + std::shared_ptr DevImpl = + sycl::detail::getSyclObjImpl(syclDevice); + std::shared_ptr CtxImpl = + sycl::detail::getSyclObjImpl(syclContext); + const sycl::detail::AdapterPtr &Adapter = CtxImpl->getAdapter(); + + ur_image_desc_t urDesc; + ur_image_format_t urFormat; + populate_ur_structs(imageDescriptor, urDesc, urFormat); + + const ur_exp_image_mem_type_t memHandleType = + (imageMemoryHandleType == image_memory_handle_type::opaque_handle) + ? ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE + : ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_USM_POINTER; + + ur_bool_t supportsSampledHandle{0}; + Adapter->call< + sycl::errc::runtime, + sycl::detail::UrApiKind::urBindlessImagesGetImageSampledHandleSupportExp>( + CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat, + memHandleType, &supportsSampledHandle); + + return supportsSampledHandle; +} + +template <> +__SYCL_EXPORT bool is_image_handle_supported( + const image_descriptor &imageDescriptor, + image_memory_handle_type imageMemoryHandleType, + const sycl::queue &syclQueue) { + return is_image_handle_supported( + imageDescriptor, imageMemoryHandleType, syclQueue.get_device(), + syclQueue.get_context()); +} + } // namespace ext::oneapi::experimental } // namespace _V1 } // namespace sycl diff --git a/sycl/test-e2e/bindless_images/3_channel_format.cpp b/sycl/test-e2e/bindless_images/3_channel_format.cpp index 7c502170bab11..62ae76adf32c3 100644 --- a/sycl/test-e2e/bindless_images/3_channel_format.cpp +++ b/sycl/test-e2e/bindless_images/3_channel_format.cpp @@ -13,6 +13,7 @@ #include #include +#include "helpers/common.hpp" #include // Uncomment to print additional test information @@ -45,14 +46,26 @@ int main() { syclexp::image_descriptor desc({width}, 3, sycl::image_channel_type::unsigned_int16); + // Verify ability to allocate the above image descriptor + if (!bindless_helpers::memoryAllocationSupported( + desc, syclexp::image_memory_handle_type::opaque_handle, q)) { + // We cannot allocate the opaque `image_mem` below + // Skip the test + if (ctxt.get_backend() == sycl::backend::ext_oneapi_cuda) { + std::cout << "CUDA doesn't support 3-channel formats. Skipping test.\n"; + } else { + std::cout << "Memory allocation unsupported. Skipping test.\n"; + } + return 0; + } + syclexp::image_mem imgMem(desc, dev, ctxt); q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc); q.wait_and_throw(); - // Some backends don't support 3-channel formats - // We still try to create the image, - // but we expect it to fail with UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT + // Backends which do not support 3-channel formats will have been skipped + // with the check above. syclexp::unsampled_image_handle imgHandle = sycl::ext::oneapi::experimental::create_image(imgMem, desc, dev, ctxt); @@ -77,14 +90,6 @@ int main() { } catch (const sycl::exception &ex) { const std::string_view errMsg(ex.what()); - if (ctxt.get_backend() == sycl::backend::ext_oneapi_cuda) { - if (errMsg.find("UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT") != - std::string::npos) { - std::cout << "CUDA doesn't support 3-channel formats, test passed." - << std::endl; - return 0; - } - } std::cerr << "Unexpected SYCL exception: " << errMsg << "\n"; return 1; } catch (...) { diff --git a/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp b/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp index e00fd26271372..81871a7aa43f4 100644 --- a/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp @@ -2,7 +2,7 @@ // REQUIRES: windows // DEFINE: %{link-flags}=%if cl_options %{ /clang:-ld3d12 /clang:-ldxgi /clang:-ldxguid %} %else %{ -ld3d12 -ldxgi -ldxguid %} -// RUN: %{build} %{link-flags} -o %t.out %if target-spir %{ -DDISABLE_UNORM_TESTS %} +// RUN: %{build} %{link-flags} -o %t.out // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out #pragma clang diagnostic ignored "-Waddress-of-temporary" @@ -696,8 +696,22 @@ template static bool runTest(DX12SYCLDevice &device, sycl::image_channel_type channelType, sycl::range globalSize, sycl::range localSize) { + + syclexp::image_descriptor syclImageDesc{globalSize, NChannels, channelType}; + + // Verify ability to allocate the above image descriptor. + // E.g. LevelZero does not support `unorm` channel types. + if (!bindless_helpers::memoryAllocationSupported( + syclImageDesc, syclexp::image_memory_handle_type::opaque_handle, + device.getSyclQueue())) { + // We cannot allocate the image memory, skip the test. + std::cout << "Memory allocation unsupported. Skipping test.\n"; + return true; + } + DX12InteropTest interopTestInstance( device, channelType, globalSize, localSize); + interopTestInstance.initDX12Resources(); interopTestInstance.callSYCLKernel(); bool validated = interopTestInstance.validateOutput(); @@ -732,10 +746,8 @@ int main() { validated &= runTest<1, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32, globalSize1, localSize1); -#ifndef DISABLE_UNORM_TESTS validated &= runTest<1, uint8_t, 4>( device, sycl::image_channel_type::unorm_int8, globalSize1, localSize1); -#endif validated &= runTest<1, float, 1>(device, sycl::image_channel_type::fp32, globalSize1, localSize1); validated &= runTest<1, sycl::half, 2>(device, sycl::image_channel_type::fp16, @@ -753,10 +765,8 @@ int main() { validated &= runTest<2, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32, globalSize2[0], {16, 16}); -#ifndef DISABLE_UNORM_TESTS validated &= runTest<2, uint8_t, 4>( device, sycl::image_channel_type::unorm_int8, globalSize2[1], {16, 8}); -#endif validated &= runTest<2, float, 1>(device, sycl::image_channel_type::fp32, globalSize2[2], {16, 8}); validated &= runTest<2, sycl::half, 2>(device, sycl::image_channel_type::fp16, @@ -777,10 +787,8 @@ int main() { validated &= runTest<3, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32, globalSize3[0], {16, 16, 1}); -#ifndef DISABLE_UNORM_TESTS validated &= runTest<3, uint8_t, 4>( device, sycl::image_channel_type::unorm_int8, globalSize3[1], {16, 8, 2}); -#endif validated &= runTest<3, float, 1>(device, sycl::image_channel_type::fp32, globalSize3[2], {16, 8, 1}); validated &= runTest<3, sycl::half, 2>(device, sycl::image_channel_type::fp16, diff --git a/sycl/test-e2e/bindless_images/helpers/common.hpp b/sycl/test-e2e/bindless_images/helpers/common.hpp index 8a5096779181f..4c4505cb14a14 100644 --- a/sycl/test-e2e/bindless_images/helpers/common.hpp +++ b/sycl/test-e2e/bindless_images/helpers/common.hpp @@ -1,6 +1,7 @@ #pragma once #include #include +#include #include template @@ -17,6 +18,17 @@ std::ostream &operator<<(std::ostream &os, namespace bindless_helpers { +namespace syclexp = sycl::ext::oneapi::experimental; + +bool memoryAllocationSupported(syclexp::image_descriptor &imgDesc, + syclexp::image_memory_handle_type memHandleType, + sycl::queue &syclQueue) { + auto supportedMemTypes = + syclexp::get_image_memory_support(imgDesc, syclQueue); + return std::find(supportedMemTypes.begin(), supportedMemTypes.end(), + memHandleType) != supportedMemTypes.end(); +} + template static void printTestName(std::string name, sycl::range globalSize, sycl::range localSize) { diff --git a/sycl/test-e2e/bindless_images/read_2D.cpp b/sycl/test-e2e/bindless_images/read_2D.cpp index af0d02cd8c3e7..b483c1117cf28 100644 --- a/sycl/test-e2e/bindless_images/read_2D.cpp +++ b/sycl/test-e2e/bindless_images/read_2D.cpp @@ -6,6 +6,7 @@ #include #include +#include "helpers/common.hpp" #include // Uncomment to print additional test information @@ -40,10 +41,37 @@ int main() { {width, height}, 4, sycl::image_channel_type::fp32); try { + + // Verify ability to allocate the above image descriptor + if (!bindless_helpers::memoryAllocationSupported( + desc, + sycl::ext::oneapi::experimental::image_memory_handle_type:: + opaque_handle, + q)) { + // We cannot allocate the opaque `image_mem` below + std::cout << "Memory allocation unsupported. Skipping test.\n"; + return 0; + } + // Extension: allocate memory on device and create the handle sycl::ext::oneapi::experimental::image_mem imgMem0(desc, dev, ctxt); sycl::ext::oneapi::experimental::image_mem imgMem1(desc, dev, ctxt); + // Extension: verify ability to create the unsampled image handles below + bool supportedUnsampledHandle = + sycl::ext::oneapi::experimental::is_image_handle_supported< + sycl::ext::oneapi::experimental::unsampled_image_handle>( + desc, + sycl::ext::oneapi::experimental::image_memory_handle_type:: + opaque_handle, + q); + if (!supportedUnsampledHandle) { + // We cannot create the unsampled handles below + std::cout << "Unsampled image handle creation unsupported. Skipping " + "test.\n"; + return 0; + } + // Extension: create the image and return the handle sycl::ext::oneapi::experimental::unsampled_image_handle imgHandle1 = sycl::ext::oneapi::experimental::create_image(imgMem0, desc, dev, ctxt); diff --git a/sycl/test-e2e/bindless_images/read_norm_types.cpp b/sycl/test-e2e/bindless_images/read_norm_types.cpp index 9eec60c61178c..a29832271127c 100644 --- a/sycl/test-e2e/bindless_images/read_norm_types.cpp +++ b/sycl/test-e2e/bindless_images/read_norm_types.cpp @@ -43,6 +43,24 @@ bool run_test(sycl::range globalSize, sycl::range localSize) { syclexp::image_descriptor descOut(globalSize, NChannels, sycl::image_channel_type::fp32); + // Verify ability to allocate the descIn descriptor. + // The LevelZero device does not support unnormalized types. + if (!bindless_helpers::memoryAllocationSupported( + descIn, syclexp::image_memory_handle_type::opaque_handle, q)) { + // The device does not support allocating opaque memory or creating images + // from descIn. Skip the test. + std::cout << "Memory allocation unsupported. Skipping test.\n"; + return true; + } + if (NDims == 2 && + !bindless_helpers::memoryAllocationSupported( + descIn, syclexp::image_memory_handle_type::usm_pointer, q)) { + // The device does not support allocating usm memory or creating images + // from descIn. Skip the test. + std::cout << "Memory allocation unsupported. Skipping test.\n"; + return true; + } + syclexp::image_mem_handle imgMemIn = syclexp::alloc_image_mem(descIn, q); syclexp::image_mem_handle imgMemOut = syclexp::alloc_image_mem(descOut, q); diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp index 8efe7ce1f4012..46b34eb6d482a 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp @@ -90,12 +90,12 @@ handles_t create_test_handles( template -bool run_sycl(sycl::range globalSize, sycl::range localSize, +bool run_sycl(sycl::queue syclQueue, sycl::range globalSize, + sycl::range localSize, InteropHandleT inputInteropMemHandle, InteropSemHandleT sycl_wait_semaphore_handle) { - sycl::device dev; - sycl::queue q(dev); - auto ctxt = q.get_context(); + auto dev = syclQueue.get_device(); + auto ctxt = syclQueue.get_context(); // Image descriptor - mapped to Vulkan image layout syclexp::image_descriptor desc(globalSize, NChannels, CType); @@ -134,13 +134,14 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, #ifdef TEST_SEMAPHORE_IMPORT // Extension: wait for imported semaphore - q.ext_oneapi_wait_external_semaphore(handles.sycl_wait_external_semaphore); + syclQueue.ext_oneapi_wait_external_semaphore( + handles.sycl_wait_external_semaphore); #endif std::vector out(numElems); try { sycl::buffer buf((VecType *)out.data(), outBufferRange); - q.submit([&](sycl::handler &cgh) { + syclQueue.submit([&](sycl::handler &cgh) { auto outAcc = buf.template get_access( cgh, outBufferRange); cgh.parallel_for( @@ -195,7 +196,7 @@ bool run_sycl(sycl::range globalSize, sycl::range localSize, } }); }); - q.wait_and_throw(); + syclQueue.wait_and_throw(); #ifdef TEST_SEMAPHORE_IMPORT syclexp::release_external_semaphore(handles.sycl_wait_external_semaphore, @@ -261,6 +262,27 @@ bool run_test(sycl::range dims, sycl::range localSize, uint32_t height = 1; uint32_t depth = 1; + sycl::queue syclQueue; + + // Skip `sycl::half` tests if fp16 is unsupported. + if constexpr (std::is_same_v) { + if (!syclQueue.get_device().has(sycl::aspect::fp16)) { + return true; + } + } + + // Verify SYCL device support for allocating/creating an image from the + // descriptor being tested. + // This test always maps to an `image_mem_handle` (opaque_handle). + syclexp::image_descriptor desc{dims, NChannels, CType}; + if (!bindless_helpers::memoryAllocationSupported( + desc, syclexp::image_memory_handle_type::opaque_handle, syclQueue)) { + // The device does not support allocating/creating the image with the given + // properties. Skip the test. + std::cout << "Memory allocation unsupported. Skipping test.\n"; + return true; + } + size_t numElems = dims[0]; VkImageType imgType = VK_IMAGE_TYPE_1D; @@ -449,7 +471,8 @@ bool run_test(sycl::range dims, sycl::range localSize, bool validated = run_sycl( - dims, localSize, input_mem_handle, sycl_wait_semaphore_handle); + syclQueue, dims, localSize, input_mem_handle, + sycl_wait_semaphore_handle); // Cleanup vkDestroyBuffer(vk_device, inputStagingBuffer, nullptr); diff --git a/sycl/test/abi/sycl_symbols_linux.dump b/sycl/test/abi/sycl_symbols_linux.dump index d27ff51778f7a..e7774fb5c0ed9 100644 --- a/sycl/test/abi/sycl_symbols_linux.dump +++ b/sycl/test/abi/sycl_symbols_linux.dump @@ -3047,6 +3047,8 @@ _ZN4sycl3_V13ext6oneapi12experimental23prepare_for_device_copyEPKvmRKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental23prepare_for_device_copyEPKvmRKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental23release_external_memoryENS3_12external_memERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental23release_external_memoryENS3_12external_memERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental24get_image_memory_supportERKNS3_16image_descriptorERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental24get_image_memory_supportERKNS3_16image_descriptorERKNS0_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 @@ -3055,6 +3057,10 @@ _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_11resource_ _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_11resource_fdEEENS3_18external_semaphoreENS3_29external_semaphore_descriptorIT_EERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21resource_win32_handleEEENS3_18external_semaphoreENS3_29external_semaphore_descriptorIT_EERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental25import_external_semaphoreINS3_21resource_win32_handleEEENS3_18external_semaphoreENS3_29external_semaphore_descriptorIT_EERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental25is_image_handle_supportedINS3_20sampled_image_handleEEEbRKNS3_16image_descriptorENS3_24image_memory_handle_typeERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental25is_image_handle_supportedINS3_20sampled_image_handleEEEbRKNS3_16image_descriptorENS3_24image_memory_handle_typeERKNS0_6deviceERKNS0_7contextE +_ZN4sycl3_V13ext6oneapi12experimental25is_image_handle_supportedINS3_22unsampled_image_handleEEEbRKNS3_16image_descriptorENS3_24image_memory_handle_typeERKNS0_5queueE +_ZN4sycl3_V13ext6oneapi12experimental25is_image_handle_supportedINS3_22unsampled_image_handleEEEbRKNS3_16image_descriptorENS3_24image_memory_handle_typeERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental25map_external_image_memoryENS3_12external_memERKNS3_16image_descriptorERKNS0_5queueE _ZN4sycl3_V13ext6oneapi12experimental25map_external_image_memoryENS3_12external_memERKNS3_16image_descriptorERKNS0_6deviceERKNS0_7contextE _ZN4sycl3_V13ext6oneapi12experimental26map_external_linear_memoryENS3_12external_memEmmRKNS0_5queueE diff --git a/sycl/test/abi/sycl_symbols_windows.dump b/sycl/test/abi/sycl_symbols_windows.dump index 845e6fdc9999a..6c9d773b157cc 100644 --- a/sycl/test/abi/sycl_symbols_windows.dump +++ b/sycl/test/abi/sycl_symbols_windows.dump @@ -280,6 +280,10 @@ ??$import_external_semaphore@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_fd@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z ??$import_external_semaphore@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVdevice@34@AEBVcontext@34@@Z ??$import_external_semaphore@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA?AUexternal_semaphore@01234@U?$external_semaphore_descriptor@Uresource_win32_handle@experimental@oneapi@ext@_V1@sycl@@@01234@AEBVqueue@34@@Z +??$is_image_handle_supported@Usampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVdevice@34@AEBVcontext@34@@Z +??$is_image_handle_supported@Usampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVqueue@34@@Z +??$is_image_handle_supported@Uunsampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVdevice@34@AEBVcontext@34@@Z +??$is_image_handle_supported@Uunsampled_image_handle@experimental@oneapi@ext@_V1@sycl@@@experimental@oneapi@ext@_V1@sycl@@YA_NAEBUimage_descriptor@01234@W4image_memory_handle_type@01234@AEBVqueue@34@@Z ??$update_nd_range@$00@node@experimental@oneapi@ext@_V1@sycl@@QEAAXV?$nd_range@$00@45@@Z ??$update_nd_range@$01@node@experimental@oneapi@ext@_V1@sycl@@QEAAXV?$nd_range@$01@45@@Z ??$update_nd_range@$02@node@experimental@oneapi@ext@_V1@sycl@@QEAAXV?$nd_range@$02@45@@Z @@ -4136,6 +4140,8 @@ ?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_memory_support@experimental@oneapi@ext@_V1@sycl@@YA?AV?$vector@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@AEBUimage_descriptor@12345@AEBVdevice@45@AEBVcontext@45@@Z +?get_image_memory_support@experimental@oneapi@ext@_V1@sycl@@YA?AV?$vector@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@V?$allocator@W4image_memory_handle_type@experimental@oneapi@ext@_V1@sycl@@@std@@@std@@AEBUimage_descriptor@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 diff --git a/unified-runtime/include/ur_api.h b/unified-runtime/include/ur_api.h index dcf05b2b066c7..811ebc65b0e00 100644 --- a/unified-runtime/include/ur_api.h +++ b/unified-runtime/include/ur_api.h @@ -461,6 +461,12 @@ typedef enum ur_function_t { UR_FUNCTION_ADAPTER_SET_LOGGER_CALLBACK = 266, /// Enumerator for ::urAdapterSetLoggerCallbackLevel UR_FUNCTION_ADAPTER_SET_LOGGER_CALLBACK_LEVEL = 267, + /// Enumerator for ::urBindlessImagesGetImageUnsampledHandleSupportExp + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP = 268, + /// Enumerator for ::urBindlessImagesGetImageSampledHandleSupportExp + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP = 269, + /// Enumerator for ::urBindlessImagesGetImageMemoryHandleTypeSupportExp + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP = 270, /// @cond UR_FUNCTION_FORCE_UINT32 = 0x7fffffff /// @endcond @@ -9707,6 +9713,19 @@ typedef enum ur_exp_external_semaphore_type_t { } ur_exp_external_semaphore_type_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Indicates the type of image backing memory handle. +typedef enum ur_exp_image_mem_type_t { + /// USM pointer to image memory + UR_EXP_IMAGE_MEM_TYPE_USM_POINTER = 0, + /// Opaque handle to image memory + UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE = 1, + /// @cond + UR_EXP_IMAGE_MEM_TYPE_FORCE_UINT32 = 0x7fffffff + /// @endcond + +} ur_exp_image_mem_type_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief File descriptor typedef struct ur_exp_file_descriptor_t { @@ -10201,6 +10220,115 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( /// [out][optional] returned query value size size_t *pPropSizeRet); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for allocating a given image backing memory handle type +/// with specific image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an unsampled image handle with specific +/// image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an sampled image handle with specific +/// image +/// properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet); + /////////////////////////////////////////////////////////////////////////////// /// @brief Retrieve individual image from mipmap /// @@ -14670,6 +14798,51 @@ typedef struct ur_bindless_images_image_get_info_exp_params_t { size_t **ppPropSizeRet; } ur_bindless_images_image_get_info_exp_params_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + const ur_image_desc_t **ppImageDesc; + const ur_image_format_t **ppImageFormat; + ur_exp_image_mem_type_t *pimageMemHandleType; + ur_bool_t **ppSupportedRet; +} ur_bindless_images_get_image_memory_handle_type_support_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + const ur_image_desc_t **ppImageDesc; + const ur_image_format_t **ppImageFormat; + ur_exp_image_mem_type_t *pimageMemHandleType; + ur_bool_t **ppSupportedRet; +} ur_bindless_images_get_image_unsampled_handle_support_exp_params_t; + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function parameters for +/// urBindlessImagesGetImageSampledHandleSupportExp +/// @details Each entry is a pointer to the parameter passed to the function; +/// allowing the callback the ability to modify the parameter's value +typedef struct + ur_bindless_images_get_image_sampled_handle_support_exp_params_t { + ur_context_handle_t *phContext; + ur_device_handle_t *phDevice; + const ur_image_desc_t **ppImageDesc; + const ur_image_format_t **ppImageFormat; + ur_exp_image_mem_type_t *pimageMemHandleType; + ur_bool_t **ppSupportedRet; +} ur_bindless_images_get_image_sampled_handle_support_exp_params_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief Function parameters for urBindlessImagesMipmapGetLevelExp /// @details Each entry is a pointer to the parameter passed to the function; diff --git a/unified-runtime/include/ur_api_funcs.def b/unified-runtime/include/ur_api_funcs.def index 33f10870bf8d5..2ba60864f5940 100644 --- a/unified-runtime/include/ur_api_funcs.def +++ b/unified-runtime/include/ur_api_funcs.def @@ -170,6 +170,9 @@ _UR_API(urBindlessImagesUnsampledImageCreateExp) _UR_API(urBindlessImagesSampledImageCreateExp) _UR_API(urBindlessImagesImageCopyExp) _UR_API(urBindlessImagesImageGetInfoExp) +_UR_API(urBindlessImagesGetImageMemoryHandleTypeSupportExp) +_UR_API(urBindlessImagesGetImageUnsampledHandleSupportExp) +_UR_API(urBindlessImagesGetImageSampledHandleSupportExp) _UR_API(urBindlessImagesMipmapGetLevelExp) _UR_API(urBindlessImagesMipmapFreeExp) _UR_API(urBindlessImagesImportExternalMemoryExp) diff --git a/unified-runtime/include/ur_ddi.h b/unified-runtime/include/ur_ddi.h index dbec33e21cbe4..553bb61a7e8c7 100644 --- a/unified-runtime/include/ur_ddi.h +++ b/unified-runtime/include/ur_ddi.h @@ -1481,6 +1481,29 @@ typedef ur_result_t(UR_APICALL *ur_pfnBindlessImagesImageGetInfoExp_t)( ur_context_handle_t, ur_exp_image_mem_native_handle_t, ur_image_info_t, void *, size_t *); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +typedef ur_result_t( + UR_APICALL *ur_pfnBindlessImagesGetImageMemoryHandleTypeSupportExp_t)( + ur_context_handle_t, ur_device_handle_t, const ur_image_desc_t *, + const ur_image_format_t *, ur_exp_image_mem_type_t, ur_bool_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +typedef ur_result_t( + UR_APICALL *ur_pfnBindlessImagesGetImageUnsampledHandleSupportExp_t)( + ur_context_handle_t, ur_device_handle_t, const ur_image_desc_t *, + const ur_image_format_t *, ur_exp_image_mem_type_t, ur_bool_t *); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Function-pointer for urBindlessImagesGetImageSampledHandleSupportExp +typedef ur_result_t( + UR_APICALL *ur_pfnBindlessImagesGetImageSampledHandleSupportExp_t)( + ur_context_handle_t, ur_device_handle_t, const ur_image_desc_t *, + const ur_image_format_t *, ur_exp_image_mem_type_t, ur_bool_t *); + /////////////////////////////////////////////////////////////////////////////// /// @brief Function-pointer for urBindlessImagesMipmapGetLevelExp typedef ur_result_t(UR_APICALL *ur_pfnBindlessImagesMipmapGetLevelExp_t)( @@ -1557,6 +1580,12 @@ typedef struct ur_bindless_images_exp_dditable_t { ur_pfnBindlessImagesSampledImageCreateExp_t pfnSampledImageCreateExp; ur_pfnBindlessImagesImageCopyExp_t pfnImageCopyExp; ur_pfnBindlessImagesImageGetInfoExp_t pfnImageGetInfoExp; + ur_pfnBindlessImagesGetImageMemoryHandleTypeSupportExp_t + pfnGetImageMemoryHandleTypeSupportExp; + ur_pfnBindlessImagesGetImageUnsampledHandleSupportExp_t + pfnGetImageUnsampledHandleSupportExp; + ur_pfnBindlessImagesGetImageSampledHandleSupportExp_t + pfnGetImageSampledHandleSupportExp; ur_pfnBindlessImagesMipmapGetLevelExp_t pfnMipmapGetLevelExp; ur_pfnBindlessImagesMipmapFreeExp_t pfnMipmapFreeExp; ur_pfnBindlessImagesImportExternalMemoryExp_t pfnImportExternalMemoryExp; diff --git a/unified-runtime/include/ur_print.h b/unified-runtime/include/ur_print.h index 533a8b3afeaa8..c2dd79cc6ba23 100644 --- a/unified-runtime/include/ur_print.h +++ b/unified-runtime/include/ur_print.h @@ -1212,6 +1212,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urPrintExpExternalSemaphoreType( enum ur_exp_external_semaphore_type_t value, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_exp_image_mem_type_t enum +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL +urPrintExpImageMemType(enum ur_exp_image_mem_type_t value, char *buffer, + const size_t buff_size, size_t *out_size); + /////////////////////////////////////////////////////////////////////////////// /// @brief Print ur_exp_file_descriptor_t struct /// @returns @@ -3065,6 +3075,46 @@ UR_APIEXPORT ur_result_t UR_APICALL urPrintBindlessImagesImageGetInfoExpParams( const struct ur_bindless_images_image_get_info_exp_params_t *params, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print +/// ur_bindless_images_get_image_memory_handle_type_support_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL +urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams( + const struct + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t + *params, + char *buffer, const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print +/// ur_bindless_images_get_image_unsampled_handle_support_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL +urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams( + const struct + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size); + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print +/// ur_bindless_images_get_image_sampled_handle_support_exp_params_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL +urPrintBindlessImagesGetImageSampledHandleSupportExpParams( + const struct + ur_bindless_images_get_image_sampled_handle_support_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size); + /////////////////////////////////////////////////////////////////////////////// /// @brief Print ur_bindless_images_mipmap_get_level_exp_params_t struct /// @returns diff --git a/unified-runtime/include/ur_print.hpp b/unified-runtime/include/ur_print.hpp index c5333f76f478e..d395f135ee771 100644 --- a/unified-runtime/include/ur_print.hpp +++ b/unified-runtime/include/ur_print.hpp @@ -529,6 +529,8 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_exp_external_mem_type_t value); inline std::ostream &operator<<(std::ostream &os, enum ur_exp_external_semaphore_type_t value); +inline std::ostream &operator<<(std::ostream &os, + enum ur_exp_image_mem_type_t value); inline std::ostream & operator<<(std::ostream &os, [[maybe_unused]] const struct ur_exp_file_descriptor_t params); @@ -1244,6 +1246,16 @@ inline std::ostream &operator<<(std::ostream &os, enum ur_function_t value) { case UR_FUNCTION_ADAPTER_SET_LOGGER_CALLBACK_LEVEL: os << "UR_FUNCTION_ADAPTER_SET_LOGGER_CALLBACK_LEVEL"; break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP: + os << "UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP"; + break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP: + os << "UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP"; + break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP: + os << "UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_" + "EXP"; + break; default: os << "unknown enumerator"; break; @@ -11186,6 +11198,25 @@ inline std::ostream &operator<<(std::ostream &os, return os; } /////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_exp_image_mem_type_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<(std::ostream &os, + enum ur_exp_image_mem_type_t value) { + switch (value) { + case UR_EXP_IMAGE_MEM_TYPE_USM_POINTER: + os << "UR_EXP_IMAGE_MEM_TYPE_USM_POINTER"; + break; + case UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE: + os << "UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE"; + break; + default: + os << "unknown enumerator"; + break; + } + return os; +} +/////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the ur_exp_file_descriptor_t type /// @returns /// std::ostream & @@ -18262,6 +18293,131 @@ inline std::ostream &operator<<( return os; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the +/// ur_bindless_images_get_image_memory_handle_type_support_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, [[maybe_unused]] const struct + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hDevice = "; + + ur::details::printPtr(os, *(params->phDevice)); + + os << ", "; + os << ".pImageDesc = "; + + ur::details::printPtr(os, *(params->ppImageDesc)); + + os << ", "; + os << ".pImageFormat = "; + + ur::details::printPtr(os, *(params->ppImageFormat)); + + os << ", "; + os << ".imageMemHandleType = "; + + os << *(params->pimageMemHandleType); + + os << ", "; + os << ".pSupportedRet = "; + + ur::details::printPtr(os, *(params->ppSupportedRet)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the +/// ur_bindless_images_get_image_unsampled_handle_support_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, [[maybe_unused]] const struct + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t + *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hDevice = "; + + ur::details::printPtr(os, *(params->phDevice)); + + os << ", "; + os << ".pImageDesc = "; + + ur::details::printPtr(os, *(params->ppImageDesc)); + + os << ", "; + os << ".pImageFormat = "; + + ur::details::printPtr(os, *(params->ppImageFormat)); + + os << ", "; + os << ".imageMemHandleType = "; + + os << *(params->pimageMemHandleType); + + os << ", "; + os << ".pSupportedRet = "; + + ur::details::printPtr(os, *(params->ppSupportedRet)); + + return os; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the +/// ur_bindless_images_get_image_sampled_handle_support_exp_params_t type +/// @returns +/// std::ostream & +inline std::ostream &operator<<( + std::ostream &os, [[maybe_unused]] const struct + ur_bindless_images_get_image_sampled_handle_support_exp_params_t *params) { + + os << ".hContext = "; + + ur::details::printPtr(os, *(params->phContext)); + + os << ", "; + os << ".hDevice = "; + + ur::details::printPtr(os, *(params->phDevice)); + + os << ", "; + os << ".pImageDesc = "; + + ur::details::printPtr(os, *(params->ppImageDesc)); + + os << ", "; + os << ".pImageFormat = "; + + ur::details::printPtr(os, *(params->ppImageFormat)); + + os << ", "; + os << ".imageMemHandleType = "; + + os << *(params->pimageMemHandleType); + + os << ", "; + os << ".pSupportedRet = "; + + ur::details::printPtr(os, *(params->ppSupportedRet)); + + return os; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the /// ur_bindless_images_mipmap_get_level_exp_params_t type @@ -21223,6 +21379,21 @@ inline ur_result_t UR_APICALL printFunctionParams(std::ostream &os, case UR_FUNCTION_BINDLESS_IMAGES_IMAGE_GET_INFO_EXP: { os << (const struct ur_bindless_images_image_get_info_exp_params_t *)params; } break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP: { + os << (const struct + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t + *)params; + } break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP: { + os << (const struct + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t *) + params; + } break; + case UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP: { + os << (const struct + ur_bindless_images_get_image_sampled_handle_support_exp_params_t *) + params; + } break; case UR_FUNCTION_BINDLESS_IMAGES_MIPMAP_GET_LEVEL_EXP: { os << (const struct ur_bindless_images_mipmap_get_level_exp_params_t *) params; diff --git a/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst b/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst index c2baba23c8db3..479afb87e4f28 100644 --- a/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst +++ b/unified-runtime/scripts/core/EXP-BINDLESS-IMAGES.rst @@ -146,11 +146,18 @@ Enums * ${X}_FUNCTION_BINDLESS_IMAGES_RELEASE_EXTERNAL_SEMAPHORE_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_WAIT_EXTERNAL_SEMAPHORE_EXP * ${X}_FUNCTION_BINDLESS_IMAGES_SIGNAL_EXTERNAL_SEMAPHORE_EXP + * ${X}_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP + * ${X}_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP + * ${X}_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP * ${x}_mem_type_t * ${X}_MEM_TYPE_IMAGE_CUBEMAP_EXP * ${X}_MEM_TYPE_IMAGE_GATHER_EXP +* {x}_exp_image_mem_type_t + * {X}_EXP_IMAGE_MEM_TYPE_USM_POINTER + * {X}_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE + Types ~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~~ * ${x}_exp_sampler_mip_properties_t @@ -182,6 +189,9 @@ Functions * ${x}BindlessImagesImageGetInfoExp * ${x}BindlessImagesMipmapGetLevelExp * ${x}BindlessImagesMipmapFreeExp + * ${x}BindlessImagesGetImageMemoryHandleTypeSupportExp + * ${x}BindlessImagesGetImageUnsampledHandleSupportExp + * ${x}BindlessImagesGetImageSampledHandleSupportExp * Interop * ${x}BindlessImagesImportExternalMemoryExp @@ -269,6 +279,13 @@ Changelog | || to DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_1D_USM_SUPPORT_EXP | | || to be more consistent with other UR enums | +----------+--------------------------------------------------------------------+ +| 22.0 || Added the following enum: | +| || - exp_image_mem_type_t | +| || Added the following APIs: | +| || - GetImageMemoryHandleTypeSupportExp | +| || - GetImageUnsampledHandleSupportExp | +| || - GetImageSampledHandleSupportExp | ++----------+-------------------------------------------------------------+ Contributors -------------------------------------------------------------------------------- diff --git a/unified-runtime/scripts/core/exp-bindless-images.yml b/unified-runtime/scripts/core/exp-bindless-images.yml index b0fb250b879b7..51e9cacc87db0 100644 --- a/unified-runtime/scripts/core/exp-bindless-images.yml +++ b/unified-runtime/scripts/core/exp-bindless-images.yml @@ -229,6 +229,16 @@ etors: - name: TIMELINE_WIN32_NT desc: "Timeline semaphore Win32 NT handle" --- #-------------------------------------------------------------------------- +type: enum +desc: "Indicates the type of image backing memory handle." +class: $xBindlessImages +name: $x_exp_image_mem_type_t +etors: + - name: USM_POINTER + desc: "USM pointer to image memory" + - name: OPAQUE_HANDLE + desc: "Opaque handle to image memory" +--- #-------------------------------------------------------------------------- type: struct desc: "File descriptor" name: $x_exp_file_descriptor_t @@ -643,6 +653,90 @@ returns: - $X_RESULT_ERROR_OUT_OF_HOST_MEMORY --- #-------------------------------------------------------------------------- type: function +desc: "Query support for allocating a given image backing memory handle type with specific image properties" +class: $xBindlessImages +name: GetImageMemoryHandleTypeSupportExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_device_handle_t + name: hDevice + desc: "[in] handle of the device object" + - type: "const $x_image_desc_t*" + name: pImageDesc + desc: "[in] pointer to image description" + - type: "const $x_image_format_t*" + name: pImageFormat + desc: "[in] pointer to image format specification" + - type: $x_exp_image_mem_type_t + name: imageMemHandleType + desc: "[in] type of image backing memory handle to query support for" + - type: $x_bool_t* + name: pSupportedRet + desc: "[out] returned indication of support for allocating the given image backing memory handle type" +returns: + - $X_RESULT_ERROR_INVALID_DEVICE + - $X_RESULT_ERROR_INVALID_CONTEXT +--- #-------------------------------------------------------------------------- +type: function +desc: "Query support for creating an unsampled image handle with specific image properties" +class: $xBindlessImages +name: GetImageUnsampledHandleSupportExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_device_handle_t + name: hDevice + desc: "[in] handle of the device object" + - type: "const $x_image_desc_t*" + name: pImageDesc + desc: "[in] pointer to image description" + - type: "const $x_image_format_t*" + name: pImageFormat + desc: "[in] pointer to image format specification" + - type: $x_exp_image_mem_type_t + name: imageMemHandleType + desc: "[in] type of image backing memory handle to query support for" + - type: $x_bool_t* + name: pSupportedRet + desc: "[out] returned indication of support for creating unsampled image handles" +returns: + - $X_RESULT_ERROR_INVALID_DEVICE + - $X_RESULT_ERROR_INVALID_CONTEXT +--- #-------------------------------------------------------------------------- +type: function +desc: "Query support for creating an sampled image handle with specific image properties" +class: $xBindlessImages +name: GetImageSampledHandleSupportExp +ordinal: "0" +params: + - type: $x_context_handle_t + name: hContext + desc: "[in] handle of the context object" + - type: $x_device_handle_t + name: hDevice + desc: "[in] handle of the device object" + - type: "const $x_image_desc_t*" + name: pImageDesc + desc: "[in] pointer to image description" + - type: "const $x_image_format_t*" + name: pImageFormat + desc: "[in] pointer to image format specification" + - type: $x_exp_image_mem_type_t + name: imageMemHandleType + desc: "[in] type of image backing memory handle to query support for" + - type: $x_bool_t* + name: pSupportedRet + desc: "[out] returned indication of support for creating sampled image handles" +returns: + - $X_RESULT_ERROR_INVALID_DEVICE + - $X_RESULT_ERROR_INVALID_CONTEXT +--- #-------------------------------------------------------------------------- +type: function desc: "Retrieve individual image from mipmap" class: $xBindlessImages name: MipmapGetLevelExp diff --git a/unified-runtime/scripts/core/registry.yml b/unified-runtime/scripts/core/registry.yml index f2e4e1f251e67..e9d030ddeb994 100644 --- a/unified-runtime/scripts/core/registry.yml +++ b/unified-runtime/scripts/core/registry.yml @@ -649,6 +649,15 @@ etors: - name: ADAPTER_SET_LOGGER_CALLBACK_LEVEL desc: Enumerator for $xAdapterSetLoggerCallbackLevel value: '267' +- name: BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP + desc: Enumerator for $xBindlessImagesGetImageUnsampledHandleSupportExp + value: '268' +- name: BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP + desc: Enumerator for $xBindlessImagesGetImageSampledHandleSupportExp + value: '269' +- name: BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP + desc: Enumerator for $xBindlessImagesGetImageMemoryHandleTypeSupportExp + value: '270' --- type: enum desc: Defines structure types diff --git a/unified-runtime/source/adapters/cuda/image.cpp b/unified-runtime/source/adapters/cuda/image.cpp index 23c7b1423108d..4adf8d491c5a3 100644 --- a/unified-runtime/source/adapters/cuda/image.cpp +++ b/unified-runtime/source/adapters/cuda/image.cpp @@ -1005,6 +1005,385 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( } } +bool verifyStandardImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType) { + // Verify standard image dimensions are within device limits. + size_t maxImageWidth, maxImageHeight, maxImageDepth; + + if (pImageDesc->depth != 0 && pImageDesc->type == UR_MEM_TYPE_IMAGE3D) { + + // Verify for standard 3D images. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_WIDTH, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_HEIGHT, + sizeof(size_t), &maxImageHeight, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_DEPTH, + sizeof(size_t), &maxImageDepth, nullptr)); + if ((pImageDesc->width > maxImageWidth) || + (pImageDesc->height > maxImageHeight) || + (pImageDesc->depth > maxImageDepth)) { + return false; + } + } else if (pImageDesc->height != 0 && pImageDesc->numMipLevel == 1 && + pImageDesc->type == UR_MEM_TYPE_IMAGE2D) { + + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + // Verify for standard 2D images backed by linear memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP, + sizeof(size_t), &maxImageHeight, nullptr)); + + size_t maxImageLinearPitch; + UR_CHECK_ERROR( + urDeviceGetInfo(hDevice, UR_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH_EXP, + sizeof(size_t), &maxImageLinearPitch, nullptr)); + if (pImageDesc->rowPitch > maxImageLinearPitch) { + return false; + } + } else { + // Verify for standard 2D images backed by opaque memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE2D_MAX_HEIGHT, + sizeof(size_t), &maxImageHeight, nullptr)); + } + + if ((pImageDesc->width > maxImageWidth) || + (pImageDesc->height > maxImageHeight)) { + return false; + } + } else if (pImageDesc->width != 0 && pImageDesc->numMipLevel == 1 && + pImageDesc->type == UR_MEM_TYPE_IMAGE1D) { + + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + // Verify for standard 1D images backed by linear memory. + // + /// TODO: We have a query for `max_image_linear_width`, however, that + /// query is for 2D textures (at least as far as the CUDA/HIP + /// implementations go). We should split the `max_image_linear_width` + /// query into 1D and 2D variants to ensure that 1D image dimensions + /// can be properly verified and used to the fullest extent. + int32_t maxImageLinearWidth; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxImageLinearWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LINEAR_WIDTH, hDevice->get())); + maxImageWidth = static_cast(maxImageLinearWidth); + } else { + // Verify for standard 1D images backed by opaque memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE, + sizeof(size_t), &maxImageWidth, nullptr)); + } + if ((pImageDesc->width > maxImageWidth)) { + return false; + } + } + + return true; +} + +bool verifyMipmapImageSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify mipmap image dimensions are within device limits. + size_t maxImageWidth, maxImageHeight; + + if (pImageDesc->height != 0 && pImageDesc->numMipLevel > 1 && + pImageDesc->type == UR_MEM_TYPE_IMAGE2D) { + // Verify for 2D mipmap images. + int32_t maxMipmapWidth, maxMipmapHeight; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxMipmapWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_WIDTH, + hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxMipmapHeight, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_MIPMAPPED_HEIGHT, + hDevice->get())); + maxImageWidth = static_cast(maxMipmapWidth); + maxImageHeight = static_cast(maxMipmapHeight); + + if ((pImageDesc->width > maxImageWidth) || + (pImageDesc->height > maxImageHeight)) { + return false; + } + } else if (pImageDesc->width != 0 && pImageDesc->numMipLevel > 1 && + pImageDesc->type == UR_MEM_TYPE_IMAGE1D) { + // Verify for 1D mipmap images. + int32_t maxMipmapWidth; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxMipmapWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_MIPMAPPED_WIDTH, + hDevice->get())); + maxImageWidth = static_cast(maxMipmapWidth); + if ((pImageDesc->width > maxImageWidth)) { + return false; + } + } + + return true; +} + +bool verifyLayeredImageSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify layered image dimensions are within device limits. + size_t maxImageWidth, maxImageHeight, maxImageLayers; + + if (pImageDesc->type == UR_MEM_TYPE_IMAGE1D_ARRAY) { + // Take the smaller of maximum surface and maximum texture width, as we do + // for `UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH` and others. + int32_t maxTextureLayeredWidth, maxSurfaceLayeredWidth; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTextureLayeredWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_WIDTH, hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfaceLayeredWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_WIDTH, hDevice->get())); + + maxImageWidth = static_cast( + std::min(maxTextureLayeredWidth, maxSurfaceLayeredWidth)); + + // Take the smaller of maximum surface and maximum texture layers, as we do + // for `UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH` and others. + int32_t maxTextureLayers, maxSurfaceLayers; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTextureLayers, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE1D_LAYERED_LAYERS, + hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfaceLayers, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE1D_LAYERED_LAYERS, + hDevice->get())); + + maxImageLayers = + static_cast(std::min(maxTextureLayers, maxSurfaceLayers)); + + if (pImageDesc->width > maxImageWidth || + pImageDesc->arraySize > maxImageLayers) { + return false; + } + + } else if (pImageDesc->type == UR_MEM_TYPE_IMAGE2D_ARRAY) { + // Take the smaller of maximum surface and maximum texture width and height, + // as we do for `UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH` and others. + int32_t maxTextureLayeredWidth, maxSurfaceLayeredWidth; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTextureLayeredWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_WIDTH, hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfaceLayeredWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_WIDTH, hDevice->get())); + + int32_t maxTextureLayeredHeight, maxSurfaceLayeredHeight; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTextureLayeredHeight, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_HEIGHT, hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfaceLayeredHeight, + CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_HEIGHT, hDevice->get())); + + maxImageWidth = static_cast( + std::min(maxTextureLayeredWidth, maxSurfaceLayeredWidth)); + + maxImageHeight = static_cast( + std::min(maxTextureLayeredHeight, maxSurfaceLayeredHeight)); + + // Take the smaller of maximum surface and maximum texture layers, as we do + // for `UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH` and others. + int32_t maxTextureLayers, maxSurfaceLayers; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTextureLayers, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_LAYERED_LAYERS, + hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfaceLayers, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACE2D_LAYERED_LAYERS, + hDevice->get())); + + maxImageLayers = static_cast( + std::min(maxTextureLayeredWidth, maxSurfaceLayeredWidth)); + + if (pImageDesc->width > maxImageWidth || + pImageDesc->height > maxImageHeight || + pImageDesc->arraySize > maxImageLayers) { + return false; + } + } + + return true; +} + +bool verifyCubemapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType) { + // Verify cubemap support and whether cubemap image dimensions are within + // device limits. + size_t maxImageWidth; + + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_CUBEMAP_EXP) { + + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + // Bindless Images do not provide support for cubemaps backed by + // USM/linear memory. + return false; + } + + if (pImageDesc->arraySize != 0) { + // Bindless Images do not provide support for layered cubemaps. + return false; + } + + // Take the smaller of maximum surface and maximum texture cubemap widths, + // as we do for `UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH` and others. + int32_t maxTexCubemapWidth, maxSurfCubemapWidth; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxTexCubemapWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_SURFACECUBEMAP_WIDTH, + hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxSurfCubemapWidth, CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURECUBEMAP_WIDTH, + hDevice->get())); + + maxImageWidth = + static_cast(std::min(maxTexCubemapWidth, maxSurfCubemapWidth)); + + // Cubemaps always have equal width and height. + if (pImageDesc->width > maxImageWidth || + pImageDesc->height > maxImageWidth) { + return false; + } + } + + return true; +} + +bool verifyGatherImageSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify gather image dimensions are within device limits. + size_t maxImageWidth, maxImageHeight; + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_GATHER_EXP) { + + // Gather images only support 2D. + if (pImageDesc->height == 0 || pImageDesc->depth > 0) { + return false; + } + + int32_t maxGatherTextureWidth, maxGatherTextureHeight; + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxGatherTextureWidth, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_WIDTH, hDevice->get())); + UR_CHECK_ERROR(cuDeviceGetAttribute( + &maxGatherTextureHeight, + CU_DEVICE_ATTRIBUTE_MAXIMUM_TEXTURE2D_GATHER_HEIGHT, hDevice->get())); + + maxImageWidth = static_cast(maxGatherTextureWidth); + maxImageHeight = static_cast(maxGatherTextureHeight); + + if (pImageDesc->width > maxImageWidth || + pImageDesc->height > maxImageHeight) { + return false; + } + } + + return true; +} + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType) { + + bool supported = true; + + supported &= + verifyStandardImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyMipmapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyLayeredImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyCubemapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyGatherImageSupport(hDevice, pImageDesc, imageMemHandleType); + + // Verify 3-channel format support. + // CUDA does not allow 3-channel formats. + if (pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_RGB || + pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_RGX) { + return false; + } + + return supported; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Currently the Bindless Images extension does not allow creation of + // unsampled image handles from non-opaque (USM) memory. + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Bindless Images do not allow creation of `unsampled_image_handle`s for + // mipmap images. + if (pImageDesc->numMipLevel > 1) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_image_mem_native_handle_t hImageMem, uint32_t mipmapLevel, diff --git a/unified-runtime/source/adapters/cuda/image.hpp b/unified-runtime/source/adapters/cuda/image.hpp index 7233d1785c5ef..686441cb79ae7 100644 --- a/unified-runtime/source/adapters/cuda/image.hpp +++ b/unified-runtime/source/adapters/cuda/image.hpp @@ -33,3 +33,28 @@ ur_result_t urTextureCreate(ur_sampler_handle_t hSampler, const CUDA_RESOURCE_DESC &ResourceDesc, const unsigned int normalized_dtype_flag, ur_exp_image_native_handle_t *phRetImage); + +bool verifyStandardImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyMipmapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCubemapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyLayeredImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyGatherImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType); diff --git a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp index 32efb2a10aad6..5159915f861a6 100644 --- a/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/cuda/ur_interface_loader.cpp @@ -352,6 +352,13 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesWaitExternalSemaphoreExp; pDdiTable->pfnSignalExternalSemaphoreExp = urBindlessImagesSignalExternalSemaphoreExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + urBindlessImagesGetImageSampledHandleSupportExp; + return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/hip/device.cpp b/unified-runtime/source/adapters/hip/device.cpp index de3d17821de56..151f4579a165a 100644 --- a/unified-runtime/source/adapters/hip/device.cpp +++ b/unified-runtime/source/adapters/hip/device.cpp @@ -247,94 +247,46 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, return ReturnValue(128u); } case UR_DEVICE_INFO_IMAGE2D_MAX_HEIGHT: { - // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); assert(TexHeight >= 0); - int SurfHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); - assert(SurfHeight >= 0); - - int Min = std::min(TexHeight, SurfHeight); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexHeight)); } case UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH: { - // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); assert(TexWidth >= 0); - int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); - assert(SurfWidth >= 0); - - int Min = std::min(TexWidth, SurfWidth); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexWidth)); } case UR_DEVICE_INFO_IMAGE3D_MAX_HEIGHT: { - // Take the smaller of maximum surface and maximum texture height. int TexHeight = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); assert(TexHeight >= 0); - int SurfHeight = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfHeight, hipDeviceAttributeMaxTexture3DHeight, hDevice->get())); - assert(SurfHeight >= 0); - - int Min = std::min(TexHeight, SurfHeight); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexHeight)); } case UR_DEVICE_INFO_IMAGE3D_MAX_WIDTH: { - // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); assert(TexWidth >= 0); - int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture3DWidth, hDevice->get())); - assert(SurfWidth >= 0); - - int Min = std::min(TexWidth, SurfWidth); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexWidth)); } case UR_DEVICE_INFO_IMAGE3D_MAX_DEPTH: { - // Take the smaller of maximum surface and maximum texture depth. int TexDepth = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); assert(TexDepth >= 0); - int SurfDepth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfDepth, hipDeviceAttributeMaxTexture3DDepth, hDevice->get())); - assert(SurfDepth >= 0); - - int Min = std::min(TexDepth, SurfDepth); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexDepth)); } case UR_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE: { - // Take the smaller of maximum surface and maximum texture width. int TexWidth = 0; UR_CHECK_ERROR(hipDeviceGetAttribute( &TexWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); assert(TexWidth >= 0); - int SurfWidth = 0; - UR_CHECK_ERROR(hipDeviceGetAttribute( - &SurfWidth, hipDeviceAttributeMaxTexture1DWidth, hDevice->get())); - assert(SurfWidth >= 0); - - int Min = std::min(TexWidth, SurfWidth); - - return ReturnValue(static_cast(Min)); + return ReturnValue(static_cast(TexWidth)); } case UR_DEVICE_INFO_IMAGE_MAX_ARRAY_SIZE: { return ReturnValue(size_t(0)); @@ -782,22 +734,31 @@ UR_APIEXPORT ur_result_t UR_APICALL urDeviceGetInfo(ur_device_handle_t hDevice, &tex_pitch_align, hipDeviceAttributeTexturePitchAlignment, hDevice->get())); assert(tex_pitch_align >= 0); - return ReturnValue(static_cast(tex_pitch_align)); + return ReturnValue(tex_pitch_align); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP: { - // Default values due to non-existent hipamd queries for linear sizes. - constexpr size_t MaxLinearWidth{1}; - return ReturnValue(MaxLinearWidth); + // No direct HIP equivalent. Use `hipDeviceAttributeMaxTexture2DWidth`. + int TexWidth = 0; + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexWidth, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); + assert(TexWidth >= 0); + return ReturnValue(static_cast(TexWidth)); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP: { - // Default values due to non-existent hipamd queries for linear sizes. - constexpr size_t MaxLinearHeight{1}; - return ReturnValue(MaxLinearHeight); + // No direct HIP equivalent. Use `hipDeviceAttributeMaxTexture2DHeight`. + int TexHeight = 0; + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexHeight, hipDeviceAttributeMaxTexture2DHeight, hDevice->get())); + assert(TexHeight >= 0); + return ReturnValue(static_cast(TexHeight)); } case UR_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH_EXP: { - // Default values due to non-existent hipamd queries for linear sizes. - constexpr size_t MaxLinearPitch{1}; - return ReturnValue(MaxLinearPitch); + // No direct HIP equivalent. Use `hipDeviceAttributeMaxTexture2DWidth`. + int TexPitch = 0; + UR_CHECK_ERROR(hipDeviceGetAttribute( + &TexPitch, hipDeviceAttributeMaxTexture2DWidth, hDevice->get())); + assert(TexPitch >= 0); + return ReturnValue(static_cast(TexPitch)); } case UR_DEVICE_INFO_MIPMAP_SUPPORT_EXP: { // HIP supports mipmaps. diff --git a/unified-runtime/source/adapters/hip/image.cpp b/unified-runtime/source/adapters/hip/image.cpp index 40d69e4cd91e1..93914582fd782 100644 --- a/unified-runtime/source/adapters/hip/image.cpp +++ b/unified-runtime/source/adapters/hip/image.cpp @@ -1007,6 +1007,237 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( #endif } +bool verifyStandardImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType) { + // Verify standard image dimensions are within device limits. + size_t maxImageWidth, maxImageHeight, maxImageDepth; + + if (pImageDesc->depth != 0 && pImageDesc->type == UR_MEM_TYPE_IMAGE3D) { + + // Verify for standard 3D images. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_WIDTH, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_HEIGHT, + sizeof(size_t), &maxImageHeight, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE3D_MAX_DEPTH, + sizeof(size_t), &maxImageDepth, nullptr)); + if ((pImageDesc->width > maxImageWidth) || + (pImageDesc->height > maxImageHeight) || + (pImageDesc->depth > maxImageDepth)) { + return false; + } + } else if (pImageDesc->height != 0 && + pImageDesc->type == UR_MEM_TYPE_IMAGE2D) { + + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + // Verify for standard 2D images backed by linear memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_MAX_IMAGE_LINEAR_WIDTH_EXP, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT_EXP, + sizeof(size_t), &maxImageHeight, nullptr)); + + size_t maxImageLinearPitch; + UR_CHECK_ERROR( + urDeviceGetInfo(hDevice, UR_DEVICE_INFO_MAX_IMAGE_LINEAR_PITCH_EXP, + sizeof(size_t), &maxImageLinearPitch, nullptr)); + if (pImageDesc->rowPitch > maxImageLinearPitch) { + return false; + } + } else { + // Verify for standard 2D images backed by opaque memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE2D_MAX_WIDTH, + sizeof(size_t), &maxImageWidth, nullptr)); + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, UR_DEVICE_INFO_IMAGE2D_MAX_HEIGHT, + sizeof(size_t), &maxImageHeight, nullptr)); + } + + if ((pImageDesc->width > maxImageWidth) || + (pImageDesc->height > maxImageHeight)) { + return false; + } + } else if (pImageDesc->width != 0 && + pImageDesc->type == UR_MEM_TYPE_IMAGE1D) { + + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + // Verify for standard 1D images backed by linear memory. + // + /// TODO: We have a query for `max_image_linear_width`, however, that + /// query is for 2D textures (at least as far as the CUDA/HIP + /// implementations go). We should split the `max_image_linear_width` + /// query into 1D and 2D variants to ensure that 1D image dimensions + /// can be properly verified and used to the fullest extent. + int32_t maxImageLinearWidth; + UR_CHECK_ERROR(hipDeviceGetAttribute(&maxImageLinearWidth, + hipDeviceAttributeMaxTexture1DLinear, + hDevice->get())); + maxImageWidth = static_cast(maxImageLinearWidth); + } else { + // Verify for standard 1D images backed by opaque memory. + UR_CHECK_ERROR(urDeviceGetInfo(hDevice, + UR_DEVICE_INFO_IMAGE_MAX_BUFFER_SIZE, + sizeof(size_t), &maxImageWidth, nullptr)); + } + if ((pImageDesc->width > maxImageWidth)) { + return false; + } + } + + return true; +} + +bool verifyMipmapImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify mipmap image support. + // Mimpaps are not currently supported for the AMD target. + if (pImageDesc->numMipLevel > 1) { + return false; + } + + return true; +} + +bool verifyCubemapImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify cubemap image support. + // Cubemaps are not currently supported for the AMD target. + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_CUBEMAP_EXP) { + return false; + } + + return true; +} + +bool verifyLayeredImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify layered image support. + // Layered images are not currently supported for the AMD target. + if ((pImageDesc->type == UR_MEM_TYPE_IMAGE1D_ARRAY) || + pImageDesc->type == UR_MEM_TYPE_IMAGE2D_ARRAY) { + return false; + } + + return true; +} + +bool verifyGatherImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify gather image support. + // Gather images are not currently supported for the AMD target. + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_GATHER_EXP) { + return false; + } + + return true; +} + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType) { + + bool supported = true; + + supported &= + verifyStandardImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyMipmapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyLayeredImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyCubemapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyGatherImageSupport(hDevice, pImageDesc, imageMemHandleType); + + // Verify 3-channel format support. + // HIP does not allow 3-channel formats. + if (pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_RGB || + pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_RGX) { + return false; + } + + return supported; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Currently Bindless Images do not allow creation of unsampled image handles + // from non-opaque (USM) memory. + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Bindless Images do not allow creation of `unsampled_image_handle`s for + // mipmap images. + if (pImageDesc->numMipLevel > 1) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_image_mem_native_handle_t hImageMem, uint32_t mipmapLevel, diff --git a/unified-runtime/source/adapters/hip/image.hpp b/unified-runtime/source/adapters/hip/image.hpp index ef299ffd1194c..a4a69e3cbbc76 100644 --- a/unified-runtime/source/adapters/hip/image.hpp +++ b/unified-runtime/source/adapters/hip/image.hpp @@ -32,3 +32,28 @@ ur_result_t urTextureCreate(ur_sampler_handle_t hSampler, const HIP_RESOURCE_DESC &ResourceDesc, const unsigned int normalized_dtype_flag, ur_exp_image_native_handle_t *phRetImage); + +bool verifyStandardImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyMipmapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCubemapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyLayeredImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyGatherImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType); diff --git a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp index b7df579c68ece..b32dcff5ba89a 100644 --- a/unified-runtime/source/adapters/hip/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/hip/ur_interface_loader.cpp @@ -349,6 +349,13 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesWaitExternalSemaphoreExp; pDdiTable->pfnSignalExternalSemaphoreExp = urBindlessImagesSignalExternalSemaphoreExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + urBindlessImagesGetImageSampledHandleSupportExp; + return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/level_zero/image.cpp b/unified-runtime/source/adapters/level_zero/image.cpp index 7770c7d2126c4..56dbf628cf037 100644 --- a/unified-runtime/source/adapters/level_zero/image.cpp +++ b/unified-runtime/source/adapters/level_zero/image.cpp @@ -15,22 +15,13 @@ #include "logger/ur_logger.hpp" #include "memory.hpp" #include "sampler.hpp" +#include "ur_api.h" #include "ur_interface_loader.hpp" #include "loader/ze_loader.h" namespace { -bool Is3ChannelOrder(ur_image_channel_order_t ChannelOrder) { - switch (ChannelOrder) { - case UR_IMAGE_CHANNEL_ORDER_RGB: - case UR_IMAGE_CHANNEL_ORDER_RGX: - return true; - default: - return false; - } -} - } // namespace namespace ur::level_zero { @@ -68,9 +59,9 @@ ur_result_t urBindlessImagesImageCopyExp( bool UseCopyEngine = hQueue->useCopyEngine(/*PreferCopyEngine*/ true); // Due to the limitation of the copy engine, disable usage of Copy Engine // Given 3 channel image - if (Is3ChannelOrder( + if (is3ChannelOrder( ur_cast(pSrcImageFormat->channelOrder)) || - Is3ChannelOrder( + is3ChannelOrder( ur_cast(pDstImageFormat->channelOrder))) { UseCopyEngine = false; } diff --git a/unified-runtime/source/adapters/level_zero/image_common.cpp b/unified-runtime/source/adapters/level_zero/image_common.cpp index d80ca7d9b8522..83b9b35ef07d3 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/image_common.cpp @@ -515,6 +515,16 @@ getImageFormatTypeAndSize(const ur_image_format_t *ImageFormat) { } // namespace +bool is3ChannelOrder(ur_image_channel_order_t ChannelOrder) { + switch (ChannelOrder) { + case UR_IMAGE_CHANNEL_ORDER_RGB: + case UR_IMAGE_CHANNEL_ORDER_RGX: + return true; + default: + return false; + } +} + /// Construct ZE image desc from UR image format and desc. ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat, const ur_image_desc_t *ImageDesc, @@ -896,6 +906,144 @@ ur_result_t bindlessImagesHandleCopyFlags( } } +bool verifyStandardImageSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + + // Verify standard image dimensions are within device limits. + if (pImageDesc->depth != 0 && pImageDesc->type == UR_MEM_TYPE_IMAGE3D) { + if ((hDevice->ZeDeviceImageProperties->maxImageDims3D == 0) || + (pImageDesc->width > + hDevice->ZeDeviceImageProperties->maxImageDims3D) || + (pImageDesc->height > + hDevice->ZeDeviceImageProperties->maxImageDims3D) || + (pImageDesc->depth > + hDevice->ZeDeviceImageProperties->maxImageDims3D)) { + return false; + } + } else if (pImageDesc->height != 0 && + pImageDesc->type == UR_MEM_TYPE_IMAGE2D) { + if (((hDevice->ZeDeviceImageProperties->maxImageDims2D == 0) || + (pImageDesc->width > + hDevice->ZeDeviceImageProperties->maxImageDims2D) || + (pImageDesc->height > + hDevice->ZeDeviceImageProperties->maxImageDims2D))) { + return false; + } + } else if (pImageDesc->width != 0 && + pImageDesc->type == UR_MEM_TYPE_IMAGE1D) { + if ((hDevice->ZeDeviceImageProperties->maxImageDims1D == 0) || + (pImageDesc->width > + hDevice->ZeDeviceImageProperties->maxImageDims1D)) { + return false; + } + } + + return true; +} + +bool verifyMipmapImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify support for mipmap images. + // LevelZero currently does not support mipmap images. + if (pImageDesc->numMipLevel > 1) { + return false; + } + + return true; +} + +bool verifyCubemapImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify support for cubemap images. + // LevelZero current does not support cubemap images. + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_CUBEMAP_EXP) { + return false; + } + + return true; +} + +bool verifyLayeredImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType) { + // Verify support for layered images. + // Bindless Images do not provide support for layered images/image arrays + // backed by USM pointers. + if (((pImageDesc->type == UR_MEM_TYPE_IMAGE1D_ARRAY) || + (pImageDesc->type == UR_MEM_TYPE_IMAGE2D_ARRAY)) && + imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + return false; + } + + return true; +} + +bool verifyGatherImageSupport( + [[maybe_unused]] const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType) { + // Verify support for gather images. + // LevelZero current does not support gather images. + if (pImageDesc->type == UR_MEM_TYPE_IMAGE_GATHER_EXP) { + return false; + } + + return true; +} + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType) { + + bool supported = true; + + supported &= + verifyStandardImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyMipmapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyLayeredImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyCubemapImageSupport(hDevice, pImageDesc, imageMemHandleType); + + supported &= + verifyGatherImageSupport(hDevice, pImageDesc, imageMemHandleType); + + // Verify 3-channel format support. + // LevelZero allows 3-channel formats for `uchar` and `ushort`. + if (is3ChannelOrder(pImageFormat->channelOrder)) { + switch (pImageFormat->channelType) { + default: + return false; + case UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT8: + case UR_IMAGE_CHANNEL_TYPE_UNSIGNED_INT16: + break; + } + } + + // Verify unnormalized channel type support. + // LevelZero currently doesn't support unnormalized channel types. + switch (pImageFormat->channelType) { + default: + break; + case UR_IMAGE_CHANNEL_TYPE_UNORM_INT8: + case UR_IMAGE_CHANNEL_TYPE_UNORM_INT16: + return false; + } + + return supported; +} + namespace ur::level_zero { ur_result_t urUSMPitchedAllocExp(ur_context_handle_t hContext, @@ -1396,4 +1544,68 @@ ur_result_t urBindlessImagesMapExternalLinearMemoryExp( return UR_RESULT_SUCCESS; } +ur_result_t urBindlessImagesGetImageMemoryHandleTypeSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + return UR_RESULT_SUCCESS; +} + +ur_result_t urBindlessImagesGetImageUnsampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Currently the Bindless Images extension does not allow creation of + // unsampled image handles from non-opaque (USM) memory. + if (imageMemHandleType == UR_EXP_IMAGE_MEM_TYPE_USM_POINTER) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Bindless Images do not allow creation of `unsampled_image_handle`s for + // mipmap images. + if (pImageDesc->numMipLevel > 1) { + *pSupportedRet = false; + return UR_RESULT_SUCCESS; + } + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + + return UR_RESULT_SUCCESS; +} + +ur_result_t urBindlessImagesGetImageSampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet) { + UR_ASSERT(std::find(hContext->getDevices().begin(), + hContext->getDevices().end(), + hDevice) != hContext->getDevices().end(), + UR_RESULT_ERROR_INVALID_CONTEXT); + + // Verify support for common image properties (dims, channel types, image + // types, etc.). + *pSupportedRet = verifyCommonImagePropertiesSupport( + hDevice, pImageDesc, pImageFormat, imageMemHandleType); + + return UR_RESULT_SUCCESS; +} + } // namespace ur::level_zero diff --git a/unified-runtime/source/adapters/level_zero/image_common.hpp b/unified-runtime/source/adapters/level_zero/image_common.hpp index 92778433c206d..8df10e528b06b 100644 --- a/unified-runtime/source/adapters/level_zero/image_common.hpp +++ b/unified-runtime/source/adapters/level_zero/image_common.hpp @@ -62,3 +62,30 @@ ur_result_t bindlessImagesHandleCopyFlags( ur_exp_image_copy_flags_t imageCopyFlags, ze_command_list_handle_t ZeCommandList, ze_event_handle_t zeSignalEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents); + +bool is3ChannelOrder(ur_image_channel_order_t ChannelOrder); + +bool verifyStandardImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyMipmapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCubemapImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyLayeredImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyGatherImageSupport(const ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, + ur_exp_image_mem_type_t imageMemHandleType); + +bool verifyCommonImagePropertiesSupport( + const ur_device_handle_t hDevice, const ur_image_desc_t *pImageDesc, + const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType); diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp index 28a1301d7254f..908d20d6d9305 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.cpp @@ -83,6 +83,12 @@ UR_APIEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnImageCopyExp = ur::level_zero::urBindlessImagesImageCopyExp; pDdiTable->pfnImageGetInfoExp = ur::level_zero::urBindlessImagesImageGetInfoExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + ur::level_zero::urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + ur::level_zero::urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + ur::level_zero::urBindlessImagesGetImageSampledHandleSupportExp; pDdiTable->pfnMipmapGetLevelExp = ur::level_zero::urBindlessImagesMipmapGetLevelExp; pDdiTable->pfnMipmapFreeExp = ur::level_zero::urBindlessImagesMipmapFreeExp; diff --git a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp index 7174fba5757fc..0213e94dc8c84 100644 --- a/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp +++ b/unified-runtime/source/adapters/level_zero/ur_interface_loader.hpp @@ -558,6 +558,18 @@ ur_result_t urBindlessImagesImageCopyExp( ur_result_t urBindlessImagesImageGetInfoExp( ur_context_handle_t hContext, ur_exp_image_mem_native_handle_t hImageMem, ur_image_info_t propName, void *pPropValue, size_t *pPropSizeRet); +ur_result_t urBindlessImagesGetImageMemoryHandleTypeSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet); +ur_result_t urBindlessImagesGetImageUnsampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet); +ur_result_t urBindlessImagesGetImageSampledHandleSupportExp( + ur_context_handle_t hContext, ur_device_handle_t hDevice, + const ur_image_desc_t *pImageDesc, const ur_image_format_t *pImageFormat, + ur_exp_image_mem_type_t imageMemHandleType, ur_bool_t *pSupportedRet); ur_result_t urBindlessImagesMipmapGetLevelExp( ur_context_handle_t hContext, ur_device_handle_t hDevice, ur_exp_image_mem_native_handle_t hImageMem, uint32_t mipmapLevel, diff --git a/unified-runtime/source/adapters/mock/ur_mockddi.cpp b/unified-runtime/source/adapters/mock/ur_mockddi.cpp index 03ee006d1da8b..6634cde2000ff 100644 --- a/unified-runtime/source/adapters/mock/ur_mockddi.cpp +++ b/unified-runtime/source/adapters/mock/ur_mockddi.cpp @@ -8553,6 +8553,186 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t params = + {&hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback( + "urBindlessImagesGetImageMemoryHandleTypeSupportExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback( + "urBindlessImagesGetImageMemoryHandleTypeSupportExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback( + "urBindlessImagesGetImageMemoryHandleTypeSupportExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t params = { + &hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback( + "urBindlessImagesGetImageUnsampledHandleSupportExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback( + "urBindlessImagesGetImageUnsampledHandleSupportExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback( + "urBindlessImagesGetImageUnsampledHandleSupportExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageSampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) try { + ur_result_t result = UR_RESULT_SUCCESS; + + ur_bindless_images_get_image_sampled_handle_support_exp_params_t params = { + &hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + + auto beforeCallback = reinterpret_cast( + mock::getCallbacks().get_before_callback( + "urBindlessImagesGetImageSampledHandleSupportExp")); + if (beforeCallback) { + result = beforeCallback(¶ms); + if (result != UR_RESULT_SUCCESS) { + return result; + } + } + + auto replaceCallback = reinterpret_cast( + mock::getCallbacks().get_replace_callback( + "urBindlessImagesGetImageSampledHandleSupportExp")); + if (replaceCallback) { + result = replaceCallback(¶ms); + } else { + + result = UR_RESULT_SUCCESS; + } + + if (result != UR_RESULT_SUCCESS) { + return result; + } + + auto afterCallback = reinterpret_cast( + mock::getCallbacks().get_after_callback( + "urBindlessImagesGetImageSampledHandleSupportExp")); + if (afterCallback) { + return afterCallback(¶ms); + } + + return result; +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesMipmapGetLevelExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( @@ -11850,6 +12030,15 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnImageGetInfoExp = driver::urBindlessImagesImageGetInfoExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + driver::urBindlessImagesGetImageMemoryHandleTypeSupportExp; + + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + driver::urBindlessImagesGetImageUnsampledHandleSupportExp; + + pDdiTable->pfnGetImageSampledHandleSupportExp = + driver::urBindlessImagesGetImageSampledHandleSupportExp; + pDdiTable->pfnMipmapGetLevelExp = driver::urBindlessImagesMipmapGetLevelExp; pDdiTable->pfnMipmapFreeExp = driver::urBindlessImagesMipmapFreeExp; diff --git a/unified-runtime/source/adapters/native_cpu/image.cpp b/unified-runtime/source/adapters/native_cpu/image.cpp index d89990ed10c9e..8ad2354927afc 100644 --- a/unified-runtime/source/adapters/native_cpu/image.cpp +++ b/unified-runtime/source/adapters/native_cpu/image.cpp @@ -97,6 +97,39 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( [[maybe_unused]] ur_context_handle_t hContext, [[maybe_unused]] ur_device_handle_t hDevice, diff --git a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp index 8543428b4f314..fda01970ea959 100644 --- a/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/native_cpu/ur_interface_loader.cpp @@ -343,6 +343,13 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesWaitExternalSemaphoreExp; pDdiTable->pfnSignalExternalSemaphoreExp = urBindlessImagesSignalExternalSemaphoreExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + urBindlessImagesGetImageSampledHandleSupportExp; + return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/adapters/opencl/image.cpp b/unified-runtime/source/adapters/opencl/image.cpp index 0c628594bb55d..3be2e3dfc4054 100644 --- a/unified-runtime/source/adapters/opencl/image.cpp +++ b/unified-runtime/source/adapters/opencl/image.cpp @@ -97,6 +97,39 @@ UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; } +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + +UR_APIEXPORT ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + [[maybe_unused]] ur_context_handle_t hContext, + [[maybe_unused]] ur_device_handle_t hDevice, + [[maybe_unused]] const ur_image_desc_t *pImageDesc, + [[maybe_unused]] const ur_image_format_t *pImageFormat, + [[maybe_unused]] ur_exp_image_mem_type_t imageMemHandleType, + [[maybe_unused]] ur_bool_t *pSupportedRet) { + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; +} + UR_APIEXPORT ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( [[maybe_unused]] ur_context_handle_t hContext, [[maybe_unused]] ur_device_handle_t hDevice, diff --git a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp index 9ef88e97f6dcb..f4cf124ac7f7d 100644 --- a/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp +++ b/unified-runtime/source/adapters/opencl/ur_interface_loader.cpp @@ -362,6 +362,12 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( urBindlessImagesWaitExternalSemaphoreExp; pDdiTable->pfnSignalExternalSemaphoreExp = urBindlessImagesSignalExternalSemaphoreExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + urBindlessImagesGetImageSampledHandleSupportExp; return UR_RESULT_SUCCESS; } diff --git a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp index eb44285c18389..77e41f75a49e6 100644 --- a/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp +++ b/unified-runtime/source/loader/layers/tracing/ur_trcddi.cpp @@ -7091,6 +7091,181 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) { + auto pfnGetImageMemoryHandleTypeSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageMemoryHandleTypeSupportExp; + + if (nullptr == pfnGetImageMemoryHandleTypeSupportExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t params = + {&hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP, + "urBindlessImagesGetImageMemoryHandleTypeSupportExp", ¶ms); + + auto &logger = getContext()->logger; + logger.info(" ---> urBindlessImagesGetImageMemoryHandleTypeSupportExp\n"); + + ur_result_t result = pfnGetImageMemoryHandleTypeSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + getContext()->notify_end( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP, + "urBindlessImagesGetImageMemoryHandleTypeSupportExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_MEMORY_HANDLE_TYPE_SUPPORT_EXP, + ¶ms); + logger.info( + " <--- urBindlessImagesGetImageMemoryHandleTypeSupportExp({}) -> " + "{};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) { + auto pfnGetImageUnsampledHandleSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageUnsampledHandleSupportExp; + + if (nullptr == pfnGetImageUnsampledHandleSupportExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t params = { + &hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP, + "urBindlessImagesGetImageUnsampledHandleSupportExp", ¶ms); + + auto &logger = getContext()->logger; + logger.info(" ---> urBindlessImagesGetImageUnsampledHandleSupportExp\n"); + + ur_result_t result = pfnGetImageUnsampledHandleSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + getContext()->notify_end( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP, + "urBindlessImagesGetImageUnsampledHandleSupportExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_UNSAMPLED_HANDLE_SUPPORT_EXP, + ¶ms); + logger.info(" <--- urBindlessImagesGetImageUnsampledHandleSupportExp({}) " + "-> {};\n", + args_str.str(), result); + } + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageSampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) { + auto pfnGetImageSampledHandleSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageSampledHandleSupportExp; + + if (nullptr == pfnGetImageSampledHandleSupportExp) + return UR_RESULT_ERROR_UNSUPPORTED_FEATURE; + + ur_bindless_images_get_image_sampled_handle_support_exp_params_t params = { + &hContext, &hDevice, &pImageDesc, &pImageFormat, + &imageMemHandleType, &pSupportedRet}; + uint64_t instance = getContext()->notify_begin( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP, + "urBindlessImagesGetImageSampledHandleSupportExp", ¶ms); + + auto &logger = getContext()->logger; + logger.info(" ---> urBindlessImagesGetImageSampledHandleSupportExp\n"); + + ur_result_t result = pfnGetImageSampledHandleSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + getContext()->notify_end( + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP, + "urBindlessImagesGetImageSampledHandleSupportExp", ¶ms, &result, + instance); + + if (logger.getLevel() <= UR_LOGGER_LEVEL_INFO) { + std::ostringstream args_str; + ur::extras::printFunctionParams( + args_str, + UR_FUNCTION_BINDLESS_IMAGES_GET_IMAGE_SAMPLED_HANDLE_SUPPORT_EXP, + ¶ms); + logger.info( + " <--- urBindlessImagesGetImageSampledHandleSupportExp({}) -> {};\n", + args_str.str(), result); + } + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesMipmapGetLevelExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( @@ -9956,6 +10131,21 @@ __urdlllocal ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnImageGetInfoExp = ur_tracing_layer::urBindlessImagesImageGetInfoExp; + dditable.pfnGetImageMemoryHandleTypeSupportExp = + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + ur_tracing_layer::urBindlessImagesGetImageMemoryHandleTypeSupportExp; + + dditable.pfnGetImageUnsampledHandleSupportExp = + pDdiTable->pfnGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + ur_tracing_layer::urBindlessImagesGetImageUnsampledHandleSupportExp; + + dditable.pfnGetImageSampledHandleSupportExp = + pDdiTable->pfnGetImageSampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + ur_tracing_layer::urBindlessImagesGetImageSampledHandleSupportExp; + dditable.pfnMipmapGetLevelExp = pDdiTable->pfnMipmapGetLevelExp; pDdiTable->pfnMipmapGetLevelExp = ur_tracing_layer::urBindlessImagesMipmapGetLevelExp; diff --git a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp index bb4f35e269613..4b5662fd46f25 100644 --- a/unified-runtime/source/loader/layers/validation/ur_valddi.cpp +++ b/unified-runtime/source/loader/layers/validation/ur_valddi.cpp @@ -7937,6 +7937,195 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) { + auto pfnGetImageMemoryHandleTypeSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageMemoryHandleTypeSupportExp; + + if (nullptr == pfnGetImageMemoryHandleTypeSupportExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == pImageDesc) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pImageFormat) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pSupportedRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType) + return UR_RESULT_ERROR_INVALID_ENUMERATION; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + getContext()->refCountContext->logInvalidReference(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hDevice)) { + getContext()->refCountContext->logInvalidReference(hDevice); + } + + ur_result_t result = pfnGetImageMemoryHandleTypeSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) { + auto pfnGetImageUnsampledHandleSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageUnsampledHandleSupportExp; + + if (nullptr == pfnGetImageUnsampledHandleSupportExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == pImageDesc) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pImageFormat) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pSupportedRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType) + return UR_RESULT_ERROR_INVALID_ENUMERATION; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + getContext()->refCountContext->logInvalidReference(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hDevice)) { + getContext()->refCountContext->logInvalidReference(hDevice); + } + + ur_result_t result = pfnGetImageUnsampledHandleSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageSampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) { + auto pfnGetImageSampledHandleSupportExp = + getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageSampledHandleSupportExp; + + if (nullptr == pfnGetImageSampledHandleSupportExp) { + return UR_RESULT_ERROR_UNINITIALIZED; + } + + if (getContext()->enableParameterValidation) { + if (NULL == pImageDesc) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pImageFormat) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == pSupportedRet) + return UR_RESULT_ERROR_INVALID_NULL_POINTER; + + if (NULL == hContext) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (NULL == hDevice) + return UR_RESULT_ERROR_INVALID_NULL_HANDLE; + + if (UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType) + return UR_RESULT_ERROR_INVALID_ENUMERATION; + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hContext)) { + getContext()->refCountContext->logInvalidReference(hContext); + } + + if (getContext()->enableLifetimeValidation && + !getContext()->refCountContext->isReferenceValid(hDevice)) { + getContext()->refCountContext->logInvalidReference(hDevice); + } + + ur_result_t result = pfnGetImageSampledHandleSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesMipmapGetLevelExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( @@ -10845,6 +11034,21 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnImageGetInfoExp = ur_validation_layer::urBindlessImagesImageGetInfoExp; + dditable.pfnGetImageMemoryHandleTypeSupportExp = + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + ur_validation_layer::urBindlessImagesGetImageMemoryHandleTypeSupportExp; + + dditable.pfnGetImageUnsampledHandleSupportExp = + pDdiTable->pfnGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + ur_validation_layer::urBindlessImagesGetImageUnsampledHandleSupportExp; + + dditable.pfnGetImageSampledHandleSupportExp = + pDdiTable->pfnGetImageSampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + ur_validation_layer::urBindlessImagesGetImageSampledHandleSupportExp; + dditable.pfnMipmapGetLevelExp = pDdiTable->pfnMipmapGetLevelExp; pDdiTable->pfnMipmapGetLevelExp = ur_validation_layer::urBindlessImagesMipmapGetLevelExp; diff --git a/unified-runtime/source/loader/loader.def.in b/unified-runtime/source/loader/loader.def.in index 7a50a1ed98e1d..52b9fc18f863e 100644 --- a/unified-runtime/source/loader/loader.def.in +++ b/unified-runtime/source/loader/loader.def.in @@ -7,6 +7,9 @@ EXPORTS urAdapterRetain urAdapterSetLoggerCallback urAdapterSetLoggerCallbackLevel + urBindlessImagesGetImageMemoryHandleTypeSupportExp + urBindlessImagesGetImageSampledHandleSupportExp + urBindlessImagesGetImageUnsampledHandleSupportExp urBindlessImagesImageAllocateExp urBindlessImagesImageCopyExp urBindlessImagesImageFreeExp @@ -187,6 +190,9 @@ EXPORTS urPrintApiVersion urPrintBaseDesc urPrintBaseProperties + urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams + urPrintBindlessImagesGetImageSampledHandleSupportExpParams + urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams urPrintBindlessImagesImageAllocateExpParams urPrintBindlessImagesImageCopyExpParams urPrintBindlessImagesImageFreeExpParams @@ -339,6 +345,7 @@ EXPORTS urPrintExpFileDescriptor urPrintExpImageCopyFlags urPrintExpImageCopyRegion + urPrintExpImageMemType urPrintExpLaunchProperty urPrintExpLaunchPropertyId urPrintExpPeerInfo diff --git a/unified-runtime/source/loader/loader.map.in b/unified-runtime/source/loader/loader.map.in index f01263532eb8f..d1a7c8d190315 100644 --- a/unified-runtime/source/loader/loader.map.in +++ b/unified-runtime/source/loader/loader.map.in @@ -7,6 +7,9 @@ urAdapterRetain; urAdapterSetLoggerCallback; urAdapterSetLoggerCallbackLevel; + urBindlessImagesGetImageMemoryHandleTypeSupportExp; + urBindlessImagesGetImageSampledHandleSupportExp; + urBindlessImagesGetImageUnsampledHandleSupportExp; urBindlessImagesImageAllocateExp; urBindlessImagesImageCopyExp; urBindlessImagesImageFreeExp; @@ -187,6 +190,9 @@ urPrintApiVersion; urPrintBaseDesc; urPrintBaseProperties; + urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams; + urPrintBindlessImagesGetImageSampledHandleSupportExpParams; + urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams; urPrintBindlessImagesImageAllocateExpParams; urPrintBindlessImagesImageCopyExpParams; urPrintBindlessImagesImageFreeExpParams; @@ -339,6 +345,7 @@ urPrintExpFileDescriptor; urPrintExpImageCopyFlags; urPrintExpImageCopyRegion; + urPrintExpImageMemType; urPrintExpLaunchProperty; urPrintExpLaunchPropertyId; urPrintExpPeerInfo; diff --git a/unified-runtime/source/loader/ur_ldrddi.cpp b/unified-runtime/source/loader/ur_ldrddi.cpp index 23fa016df8c4a..8ba14d7b8ad1c 100644 --- a/unified-runtime/source/loader/ur_ldrddi.cpp +++ b/unified-runtime/source/loader/ur_ldrddi.cpp @@ -7147,6 +7147,135 @@ __urdlllocal ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageMemoryHandleTypeSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + + [[maybe_unused]] auto context = getContext(); + + // extract platform's function pointer table + auto dditable = reinterpret_cast(hContext)->dditable; + auto pfnGetImageMemoryHandleTypeSupportExp = + dditable->ur.BindlessImagesExp.pfnGetImageMemoryHandleTypeSupportExp; + if (nullptr == pfnGetImageMemoryHandleTypeSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // convert loader handle to platform handle + hContext = reinterpret_cast(hContext)->handle; + + // convert loader handle to platform handle + hDevice = reinterpret_cast(hDevice)->handle; + + // forward to device-platform + result = pfnGetImageMemoryHandleTypeSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageUnsampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + + [[maybe_unused]] auto context = getContext(); + + // extract platform's function pointer table + auto dditable = reinterpret_cast(hContext)->dditable; + auto pfnGetImageUnsampledHandleSupportExp = + dditable->ur.BindlessImagesExp.pfnGetImageUnsampledHandleSupportExp; + if (nullptr == pfnGetImageUnsampledHandleSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // convert loader handle to platform handle + hContext = reinterpret_cast(hContext)->handle; + + // convert loader handle to platform handle + hDevice = reinterpret_cast(hDevice)->handle; + + // forward to device-platform + result = pfnGetImageUnsampledHandleSupportExp( + hContext, hDevice, pImageDesc, pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Intercept function for +/// urBindlessImagesGetImageSampledHandleSupportExp +__urdlllocal ur_result_t UR_APICALL +urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + + [[maybe_unused]] auto context = getContext(); + + // extract platform's function pointer table + auto dditable = reinterpret_cast(hContext)->dditable; + auto pfnGetImageSampledHandleSupportExp = + dditable->ur.BindlessImagesExp.pfnGetImageSampledHandleSupportExp; + if (nullptr == pfnGetImageSampledHandleSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + // convert loader handle to platform handle + hContext = reinterpret_cast(hContext)->handle; + + // convert loader handle to platform handle + hDevice = reinterpret_cast(hDevice)->handle; + + // forward to device-platform + result = pfnGetImageSampledHandleSupportExp(hContext, hDevice, pImageDesc, + pImageFormat, imageMemHandleType, + pSupportedRet); + + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Intercept function for urBindlessImagesMipmapGetLevelExp __urdlllocal ur_result_t UR_APICALL urBindlessImagesMipmapGetLevelExp( @@ -10148,6 +10277,12 @@ UR_DLLEXPORT ur_result_t UR_APICALL urGetBindlessImagesExpProcAddrTable( pDdiTable->pfnImageCopyExp = ur_loader::urBindlessImagesImageCopyExp; pDdiTable->pfnImageGetInfoExp = ur_loader::urBindlessImagesImageGetInfoExp; + pDdiTable->pfnGetImageMemoryHandleTypeSupportExp = + ur_loader::urBindlessImagesGetImageMemoryHandleTypeSupportExp; + pDdiTable->pfnGetImageUnsampledHandleSupportExp = + ur_loader::urBindlessImagesGetImageUnsampledHandleSupportExp; + pDdiTable->pfnGetImageSampledHandleSupportExp = + ur_loader::urBindlessImagesGetImageSampledHandleSupportExp; pDdiTable->pfnMipmapGetLevelExp = ur_loader::urBindlessImagesMipmapGetLevelExp; pDdiTable->pfnMipmapFreeExp = ur_loader::urBindlessImagesMipmapFreeExp; diff --git a/unified-runtime/source/loader/ur_libapi.cpp b/unified-runtime/source/loader/ur_libapi.cpp index 41d52b781fdf5..4000e7f483dc0 100644 --- a/unified-runtime/source/loader/ur_libapi.cpp +++ b/unified-runtime/source/loader/ur_libapi.cpp @@ -7747,6 +7747,148 @@ ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return exceptionToResult(std::current_exception()); } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for allocating a given image backing memory handle type +/// with specific image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) try { + auto pfnGetImageMemoryHandleTypeSupportExp = + ur_lib::getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageMemoryHandleTypeSupportExp; + if (nullptr == pfnGetImageMemoryHandleTypeSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetImageMemoryHandleTypeSupportExp(hContext, hDevice, pImageDesc, + pImageFormat, imageMemHandleType, + pSupportedRet); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an unsampled image handle with specific +/// image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) try { + auto pfnGetImageUnsampledHandleSupportExp = + ur_lib::getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageUnsampledHandleSupportExp; + if (nullptr == pfnGetImageUnsampledHandleSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetImageUnsampledHandleSupportExp(hContext, hDevice, pImageDesc, + pImageFormat, imageMemHandleType, + pSupportedRet); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an sampled image handle with specific +/// image +/// properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) try { + auto pfnGetImageSampledHandleSupportExp = + ur_lib::getContext() + ->urDdiTable.BindlessImagesExp.pfnGetImageSampledHandleSupportExp; + if (nullptr == pfnGetImageSampledHandleSupportExp) + return UR_RESULT_ERROR_UNINITIALIZED; + + return pfnGetImageSampledHandleSupportExp(hContext, hDevice, pImageDesc, + pImageFormat, imageMemHandleType, + pSupportedRet); +} catch (...) { + return exceptionToResult(std::current_exception()); +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Retrieve individual image from mipmap /// diff --git a/unified-runtime/source/loader/ur_print.cpp b/unified-runtime/source/loader/ur_print.cpp index bf64377f0255e..eabc4327735a5 100644 --- a/unified-runtime/source/loader/ur_print.cpp +++ b/unified-runtime/source/loader/ur_print.cpp @@ -972,6 +972,14 @@ urPrintExpExternalSemaphoreType(enum ur_exp_external_semaphore_type_t value, return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t urPrintExpImageMemType(enum ur_exp_image_mem_type_t value, + char *buffer, const size_t buff_size, + size_t *out_size) { + std::stringstream ss; + ss << value; + return str_copy(&ss, buffer, buff_size, out_size); +} + ur_result_t urPrintExpFileDescriptor(const struct ur_exp_file_descriptor_t params, char *buffer, const size_t buff_size, @@ -1287,6 +1295,34 @@ ur_result_t urPrintBindlessImagesImageGetInfoExpParams( return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t urPrintBindlessImagesGetImageMemoryHandleTypeSupportExpParams( + const struct + ur_bindless_images_get_image_memory_handle_type_support_exp_params_t + *params, + char *buffer, const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + +ur_result_t urPrintBindlessImagesGetImageUnsampledHandleSupportExpParams( + const struct + ur_bindless_images_get_image_unsampled_handle_support_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + +ur_result_t urPrintBindlessImagesGetImageSampledHandleSupportExpParams( + const struct + ur_bindless_images_get_image_sampled_handle_support_exp_params_t *params, + char *buffer, const size_t buff_size, size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + ur_result_t urPrintBindlessImagesMipmapGetLevelExpParams( const struct ur_bindless_images_mipmap_get_level_exp_params_t *params, char *buffer, const size_t buff_size, size_t *out_size) { diff --git a/unified-runtime/source/ur_api.cpp b/unified-runtime/source/ur_api.cpp index 5c648d51a91c7..37c81a3b978c7 100644 --- a/unified-runtime/source/ur_api.cpp +++ b/unified-runtime/source/ur_api.cpp @@ -6779,6 +6779,121 @@ ur_result_t UR_APICALL urBindlessImagesImageGetInfoExp( return result; } +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for allocating a given image backing memory handle type +/// with specific image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageMemoryHandleTypeSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for allocating the given image + /// backing memory handle type + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an unsampled image handle with specific +/// image properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageUnsampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating unsampled image + /// handles + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + +/////////////////////////////////////////////////////////////////////////////// +/// @brief Query support for creating an sampled image handle with specific +/// image +/// properties +/// +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_UNINITIALIZED +/// - ::UR_RESULT_ERROR_DEVICE_LOST +/// - ::UR_RESULT_ERROR_ADAPTER_SPECIFIC +/// - ::UR_RESULT_ERROR_INVALID_NULL_HANDLE +/// + `NULL == hContext` +/// + `NULL == hDevice` +/// - ::UR_RESULT_ERROR_INVALID_NULL_POINTER +/// + `NULL == pImageDesc` +/// + `NULL == pImageFormat` +/// + `NULL == pSupportedRet` +/// - ::UR_RESULT_ERROR_INVALID_ENUMERATION +/// + `::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE < imageMemHandleType` +/// - ::UR_RESULT_ERROR_INVALID_DEVICE +/// - ::UR_RESULT_ERROR_INVALID_CONTEXT +ur_result_t UR_APICALL urBindlessImagesGetImageSampledHandleSupportExp( + /// [in] handle of the context object + ur_context_handle_t hContext, + /// [in] handle of the device object + ur_device_handle_t hDevice, + /// [in] pointer to image description + const ur_image_desc_t *pImageDesc, + /// [in] pointer to image format specification + const ur_image_format_t *pImageFormat, + /// [in] type of image backing memory handle to query support for + ur_exp_image_mem_type_t imageMemHandleType, + /// [out] returned indication of support for creating sampled image + /// handles + ur_bool_t *pSupportedRet) { + ur_result_t result = UR_RESULT_SUCCESS; + return result; +} + /////////////////////////////////////////////////////////////////////////////// /// @brief Retrieve individual image from mipmap ///