[SYCL][Bindless][UR][E2E] Add image memory and handle support queries (#17865)

This patch introduces two new SYCL queries for Bindless Images.

`get_image_memory_support` returns a vector of supported image backing
memory types for a device given an `image_descriptor.

`is_image_handle_supported` returns a boolean indicating whether the
device supports creation of either a `sampled_image_handle` or
`unsampled_image_handle`, given an `image_descriptor` and image backing
memory type.

Some tests are updated to use these queries to filter out unsupported
image properties, e.g. allocating `unorm` channel types on the LevelZero
backend.

Additionally some fixes are made to the UR HIP device queries that
pertain to supported image properties.

The HIP queries for `UR_DEVICE_INFO_IMAGE<N>D_MAX_<WIDTH/HEIGHT/DEPTH>`
have been amended to remove redundant calls to `hipDeviceGetAttribute`
(as this was being called twice unnecessarily, with the same query
parameter).

The HIP queries for
`UR_DEVICE_INFO_MAX_IMAGE_LINEAR_<WIDTH/HEIGHT/PITCH>_EXP` have been
amended to return plausible values for the HIP backend (instead of
returning `1` for each of these).

Author: przemektmalon

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<image_memory_handle_type> 
+get_image_memory_support(const image_descriptor &imageDescriptor,
+                         const sycl::device &syclDevice, 
+                         const sycl::context &syclContext);
+
+std::vector<image_memory_handle_type> 
+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 <typename ImageHandleType>
+bool
+is_image_handle_supported(const image_descriptor &imageDescriptor,
+                          image_memory_handle_type imageMemoryHandleType,
+                          const sycl::device &syclDevice,
+                          const sycl::context &syclContext);
+
+template <typename ImageHandleType>
+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:

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<image_memory_handle_type>
+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<image_memory_handle_type>
+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 <typename ImageHandleType>
+__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 <typename ImageHandleType>
+__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

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

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<image_memory_handle_type>
+get_image_memory_support(const image_descriptor &imageDescriptor,
+                         const sycl::device &syclDevice,
+                         const sycl::context &syclContext) {
+  std::shared_ptr<sycl::detail::device_impl> DevImpl =
+      sycl::detail::getSyclObjImpl(syclDevice);
+  std::shared_ptr<sycl::detail::context_impl> 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<sycl::errc::runtime,
+                sycl::detail::UrApiKind::
+                    urBindlessImagesGetImageMemoryHandleTypeSupportExp>(
+      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<sycl::errc::runtime,
+                sycl::detail::UrApiKind::
+                    urBindlessImagesGetImageMemoryHandleTypeSupportExp>(
+      CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat,
+      ur_exp_image_mem_type_t::UR_EXP_IMAGE_MEM_TYPE_OPAQUE_HANDLE,
+      &supportsOpaqueAllocation);
+
+  std::vector<image_memory_handle_type> 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<image_memory_handle_type>
+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<unsampled_image_handle>(
+    const image_descriptor &imageDescriptor,
+    image_memory_handle_type imageMemoryHandleType,
+    const sycl::device &syclDevice, const sycl::context &syclContext) {
+  std::shared_ptr<sycl::detail::device_impl> DevImpl =
+      sycl::detail::getSyclObjImpl(syclDevice);
+  std::shared_ptr<sycl::detail::context_impl> 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<sycl::errc::runtime,
+                sycl::detail::UrApiKind::
+                    urBindlessImagesGetImageUnsampledHandleSupportExp>(
+      CtxImpl->getHandleRef(), DevImpl->getHandleRef(), &urDesc, &urFormat,
+      memHandleType, &supportsUnsampledHandle);
+
+  return supportsUnsampledHandle;
+}
+
+template <>
+__SYCL_EXPORT bool is_image_handle_supported<unsampled_image_handle>(
+    const image_descriptor &imageDescriptor,
+    image_memory_handle_type imageMemoryHandleType,
+    const sycl::queue &syclQueue) {
+  return is_image_handle_supported<unsampled_image_handle>(
+      imageDescriptor, imageMemoryHandleType, syclQueue.get_device(),
+      syclQueue.get_context());
+}
+
+template <>
+__SYCL_EXPORT bool is_image_handle_supported<sampled_image_handle>(
+    const image_descriptor &imageDescriptor,
+    image_memory_handle_type imageMemoryHandleType,
+    const sycl::device &syclDevice, const sycl::context &syclContext) {
+  std::shared_ptr<sycl::detail::device_impl> DevImpl =
+      sycl::detail::getSyclObjImpl(syclDevice);
+  std::shared_ptr<sycl::detail::context_impl> 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<sampled_image_handle>(
+    const image_descriptor &imageDescriptor,
+    image_memory_handle_type imageMemoryHandleType,
+    const sycl::queue &syclQueue) {
+  return is_image_handle_supported<sampled_image_handle>(
+      imageDescriptor, imageMemoryHandleType, syclQueue.get_device(),
+      syclQueue.get_context());
+}
+
 } // namespace ext::oneapi::experimental
 } // namespace _V1
 } // namespace sycl

sycl/test-e2e/bindless_images/3_channel_format.cpp

@@ -13,6 +13,7 @@
 #include <iostream>
 #include <sycl/detail/core.hpp>
 
+#include "helpers/common.hpp"
 #include <sycl/ext/oneapi/bindless_images.hpp>
 
 // 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 (...) {