bindless_images.hpp

Go to the documentation of this file.

//==----------- bindless_images.hpp --- SYCL bindless images ---------------==//
//
// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
// See https://llvm.org/LICENSE.txt for license information.
// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
//
//===----------------------------------------------------------------------===//
 
#pragma once
 
#include <sycl/context.hpp>                               // for context
#include <sycl/detail/export.hpp>                         // for __SYCL_EXPORT
#include <sycl/detail/pi.h>                               // for pi_uint64
#include <sycl/device.hpp>                                // for device
#include <sycl/ext/oneapi/bindless_images_descriptor.hpp> // for image_desc...
#include <sycl/ext/oneapi/bindless_images_interop.hpp>    // for interop_me...
#include <sycl/ext/oneapi/bindless_images_memory.hpp>     // for image_mem_...
#include <sycl/ext/oneapi/bindless_images_sampler.hpp>    // for bindless_i...
#include <sycl/image.hpp>                                 // for image_chan...
#include <sycl/queue.hpp>                                 // for queue
#include <sycl/range.hpp>                                 // for range
 
#include <assert.h>    // for assert
#include <stddef.h>    // for size_t
#include <type_traits> // for is_scalar
 
#ifdef __SYCL_DEVICE_ONLY__
#include <sycl/detail/image_ocl_types.hpp> // for __invoke__*
#endif
 
namespace sycl {
inline namespace _V1 {
namespace ext::oneapi::experimental {
 
struct unsampled_image_handle {
  using raw_image_handle_type = pi_uint64;
 
  unsampled_image_handle() : raw_handle(~0) {}
 
  unsampled_image_handle(raw_image_handle_type raw_image_handle)
      : raw_handle(raw_image_handle) {}
 
  raw_image_handle_type raw_handle;
};
 
struct sampled_image_handle {
  using raw_image_handle_type = pi_uint64;
 
  sampled_image_handle() : raw_handle(~0) {}
 
  sampled_image_handle(raw_image_handle_type raw_image_handle)
      : raw_handle(raw_image_handle) {}
 
  raw_image_handle_type raw_handle;
};
 
__SYCL_EXPORT image_mem_handle
alloc_image_mem(const image_descriptor &desc, const sycl::device &syclDevice,
                const sycl::context &syclContext);
 
__SYCL_EXPORT image_mem_handle alloc_image_mem(const image_descriptor &desc,
                                               const sycl::queue &syclQueue);
 
__SYCL_EXPORT_DEPRECATED("Distinct image frees are deprecated. "
                         "Instead use overload that accepts image_type.")
void free_image_mem(image_mem_handle handle, const sycl::device &syclDevice,
                    const sycl::context &syclContext);
 
__SYCL_EXPORT_DEPRECATED("Distinct image frees are deprecated. "
                         "Instead use overload that accepts image_type.")
void free_image_mem(image_mem_handle handle, const sycl::queue &syclQueue);
 
__SYCL_EXPORT void free_image_mem(image_mem_handle handle, image_type imageType,
                                  const sycl::device &syclDevice,
                                  const sycl::context &syclContext);
 
__SYCL_EXPORT void free_image_mem(image_mem_handle handle, image_type imageType,
                                  const sycl::queue &syclQueue);
 
__SYCL_EXPORT_DEPRECATED("Distinct mipmap allocs are deprecated. "
                         "Instead use alloc_image_mem().")
image_mem_handle alloc_mipmap_mem(const image_descriptor &desc,
                                  const sycl::device &syclDevice,
                                  const sycl::context &syclContext);
 
__SYCL_EXPORT_DEPRECATED("Distinct mipmap allocs are deprecated. "
                         "Instead use alloc_image_mem().")
image_mem_handle alloc_mipmap_mem(const image_descriptor &desc,
                                  const sycl::device &syclQueue);
 
__SYCL_EXPORT_DEPRECATED(
    "Distinct mipmap frees are deprecated. "
    "Instead use free_image_mem() that accepts image_type.")
void free_mipmap_mem(image_mem_handle handle, const sycl::device &syclDevice,
                     const sycl::context &syclContext);
 
__SYCL_EXPORT_DEPRECATED(
    "Distinct mipmap frees are deprecated. "
    "Instead use free_image_mem() that accepts image_type.")
void free_mipmap_mem(image_mem_handle handle, const sycl::queue &syclQueue);
 
__SYCL_EXPORT image_mem_handle get_mip_level_mem_handle(
    const image_mem_handle mipMem, const unsigned int level,
    const sycl::device &syclDevice, const sycl::context &syclContext);
 
__SYCL_EXPORT image_mem_handle get_mip_level_mem_handle(
    const image_mem_handle mipMem, const unsigned int level,
    const sycl::queue &syclQueue);
 
template <typename ExternalMemHandleType>
__SYCL_EXPORT interop_mem_handle import_external_memory(
    external_mem_descriptor<ExternalMemHandleType> externalMem,
    const sycl::device &syclDevice, const sycl::context &syclContext);
 
template <typename ExternalMemHandleType>
__SYCL_EXPORT interop_mem_handle import_external_memory(
    external_mem_descriptor<ExternalMemHandleType> externalMem,
    const sycl::queue &syclQueue);
 
__SYCL_EXPORT_DEPRECATED("map_external_memory_array is deprecated."
                         "use map_external_image_memory")
image_mem_handle map_external_memory_array(interop_mem_handle memHandle,
                                           const image_descriptor &desc,
                                           const sycl::device &syclDevice,
                                           const sycl::context &syclContext);
 
__SYCL_EXPORT_DEPRECATED("map_external_memory_array is deprecated."
                         "use map_external_image_memory")
image_mem_handle map_external_memory_array(interop_mem_handle memHandle,
                                           const image_descriptor &desc,
                                           const sycl::queue &syclQueue);
 
__SYCL_EXPORT
image_mem_handle map_external_image_memory(interop_mem_handle memHandle,
                                           const image_descriptor &desc,
                                           const sycl::device &syclDevice,
                                           const sycl::context &syclContext);
 
__SYCL_EXPORT
image_mem_handle map_external_image_memory(interop_mem_handle memHandle,
                                           const image_descriptor &desc,
                                           const sycl::queue &syclQueue);
 
template <typename ExternalSemaphoreHandleType>
__SYCL_EXPORT interop_semaphore_handle import_external_semaphore(
    external_semaphore_descriptor<ExternalSemaphoreHandleType>
        externalSemaphoreDesc,
    const sycl::device &syclDevice, const sycl::context &syclContext);
 
template <typename ExternalSemaphoreHandleType>
__SYCL_EXPORT interop_semaphore_handle import_external_semaphore(
    external_semaphore_descriptor<ExternalSemaphoreHandleType>
        externalSemaphoreDesc,
    const sycl::queue &syclQueue);
 
__SYCL_EXPORT void
destroy_external_semaphore(interop_semaphore_handle semaphoreHandle,
                           const sycl::device &syclDevice,
                           const sycl::context &syclContext);
 
__SYCL_EXPORT void
destroy_external_semaphore(interop_semaphore_handle semaphoreHandle,
                           const sycl::queue &syclQueue);
 
__SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle,
                                           const sycl::device &syclDevice,
                                           const sycl::context &syclContext);
 
__SYCL_EXPORT void release_external_memory(interop_mem_handle interopHandle,
                                           const sycl::queue &syclQueue);
 
__SYCL_EXPORT unsampled_image_handle
create_image(image_mem &memHandle, const image_descriptor &desc,
             const sycl::device &syclDevice, const sycl::context &syclContext);
 
__SYCL_EXPORT unsampled_image_handle create_image(image_mem &memHandle,
                                                  const image_descriptor &desc,
                                                  const sycl::queue &syclQueue);
 
__SYCL_EXPORT unsampled_image_handle
create_image(image_mem_handle memHandle, const image_descriptor &desc,
             const sycl::device &syclDevice, const sycl::context &syclContext);
 
__SYCL_EXPORT unsampled_image_handle create_image(image_mem_handle memHandle,
                                                  const image_descriptor &desc,
                                                  const sycl::queue &syclQueue);
 
__SYCL_EXPORT sampled_image_handle
create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::device &syclDevice,
             const sycl::context &syclContext);
 
__SYCL_EXPORT sampled_image_handle
create_image(void *imgMem, size_t pitch, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::queue &syclQueue);
 
__SYCL_EXPORT sampled_image_handle
create_image(image_mem &memHandle, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::device &syclDevice,
             const sycl::context &syclContext);
 
__SYCL_EXPORT sampled_image_handle
create_image(image_mem &memHandle, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::queue &syclQueue);
 
__SYCL_EXPORT sampled_image_handle
create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::device &syclDevice,
             const sycl::context &syclContext);
 
__SYCL_EXPORT sampled_image_handle
create_image(image_mem_handle memHandle, const bindless_image_sampler &sampler,
             const image_descriptor &desc, const sycl::queue &syclQueue);
 
__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle,
                                        const sycl::device &syclDevice,
                                        const sycl::context &syclContext);
 
__SYCL_EXPORT void destroy_image_handle(unsampled_image_handle &imageHandle,
                                        const sycl::queue &syclQueue);
 
__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle,
                                        const sycl::device &syclDevice,
                                        const sycl::context &syclContext);
 
__SYCL_EXPORT void destroy_image_handle(sampled_image_handle &imageHandle,
                                        const sycl::queue &syclQueue);
 
__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
                                         size_t widthInBytes, size_t height,
                                         unsigned int elementSizeBytes,
                                         const sycl::queue &syclQueue);
 
__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
                                         size_t widthInBytes, size_t height,
                                         unsigned int elementSizeBytes,
                                         const sycl::device &syclDevice,
                                         const sycl::context &syclContext);
 
__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
                                         const image_descriptor &desc,
                                         const sycl::queue &syclQueue);
 
__SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
                                         const image_descriptor &desc,
                                         const sycl::device &syclDevice,
                                         const sycl::context &syclContext);
 
__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle,
                                             const sycl::device &syclDevice,
                                             const sycl::context &syclContext);
 
__SYCL_EXPORT sycl::range<3> get_image_range(const image_mem_handle memHandle,
                                             const sycl::queue &syclQueue);
 
__SYCL_EXPORT sycl::image_channel_type
get_image_channel_type(const image_mem_handle memHandle,
                       const sycl::device &syclDevice,
                       const sycl::context &syclContext);
 
__SYCL_EXPORT sycl::image_channel_type
get_image_channel_type(const image_mem_handle memHandle,
                       const sycl::queue &syclQueue);
 
__SYCL_EXPORT unsigned int
get_image_num_channels(const image_mem_handle memHandle,
                       const sycl::device &syclDevice,
                       const sycl::context &syclContext);
 
__SYCL_EXPORT unsigned int
get_image_num_channels(const image_mem_handle memHandle,
                       const sycl::queue &syclQueue);
 
namespace detail {
 
// is sycl::vec
template <typename T> struct is_vec { static constexpr bool value = false; };
template <typename T, int N> struct is_vec<sycl::vec<T, N>> {
  static constexpr bool value = true;
};
template <typename T> inline constexpr bool is_vec_v = is_vec<T>::value;
 
// Get the number of coordinates
template <typename CoordT> constexpr size_t coord_size() {
  if constexpr (std::is_scalar_v<CoordT>) {
    return 1;
  } else {
    return CoordT::size();
  }
}
 
// bit_cast Color to a type the backend is known to accept
template <typename DataT> constexpr auto convert_color(DataT Color) {
  constexpr size_t dataSize = sizeof(DataT);
  static_assert(
      dataSize == 1 || dataSize == 2 || dataSize == 4 || dataSize == 8 ||
          dataSize == 16,
      "Expected input data type to be of size 1, 2, 4, 8, or 16 bytes.");
 
  if constexpr (dataSize == 1) {
    return sycl::bit_cast<uint8_t>(Color);
  } else if constexpr (dataSize == 2) {
    return sycl::bit_cast<uint16_t>(Color);
  } else if constexpr (dataSize == 4) {
    return sycl::bit_cast<uint32_t>(Color);
  } else if constexpr (dataSize == 8) {
    return sycl::bit_cast<sycl::vec<uint32_t, 2>>(Color);
  } else { // dataSize == 16
    return sycl::bit_cast<sycl::vec<uint32_t, 4>>(Color);
  }
}
 
// assert coords or elements of coords is of an integer type
template <typename CoordT> constexpr void assert_unsampled_coords() {
  if constexpr (std::is_scalar_v<CoordT>) {
    static_assert(std::is_same_v<CoordT, int>,
                  "Expected integer coordinate data type");
  } else {
    static_assert(is_vec_v<CoordT>, "Expected sycl::vec coordinates");
    static_assert(std::is_same_v<typename CoordT::element_type, int>,
                  "Expected integer coordinates data type");
  }
}
 
// assert coords or elements of coords is of a float type
template <typename CoordT> constexpr void assert_sampled_coords() {
  if constexpr (std::is_scalar_v<CoordT>) {
    static_assert(std::is_same_v<CoordT, float>,
                  "Expected float coordinate data type");
  } else {
    static_assert(is_vec_v<CoordT>, "Expected sycl::vec coordinates");
    static_assert(std::is_same_v<typename CoordT::element_type, float>,
                  "Expected float coordinates data type");
  }
}
 
template <typename DataT> constexpr bool is_data_size_valid() {
  return (sizeof(DataT) == 1) || (sizeof(DataT) == 2) || (sizeof(DataT) == 4) ||
         (sizeof(DataT) == 8) || (sizeof(DataT) == 16);
}
 
template <typename DataT> constexpr bool is_recognized_standard_type() {
  return is_data_size_valid<DataT>() &&
         (is_vec_v<DataT> || std::is_scalar_v<DataT> ||
          std::is_floating_point_v<DataT> || std::is_same_v<DataT, sycl::half>);
}
 
} // namespace detail
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_image for standard unsampled images is deprecated. "
                  "Instead use fetch_image.")
DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                 const CoordT &coords [[maybe_unused]]) {
  return fetch_image(imageHandle, coords);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]]) {
  detail::assert_unsampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
                "Expected input coordinate to be have 1, 2, or 3 components "
                "for 1D, 2D and 3D images, respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to read a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(
        __invoke__ImageRead<HintT>(imageHandle.raw_handle, coords));
  }
#else
  assert(false); // Bindless images not yet implemented on host
#endif
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_image for standard sampled images is deprecated. "
                  "Instead use sample_image.")
DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 const CoordT &coords [[maybe_unused]]) {
  return sample_image(imageHandle, coords);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                   const CoordT &coords [[maybe_unused]]) {
  detail::assert_sampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
                "Expected input coordinate to be have 1, 2, or 3 components "
                "for 1D, 2D and 3D images, respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to read a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(
        __invoke__ImageRead<HintT>(imageHandle.raw_handle, coords));
  }
#else
  assert(false); // Bindless images not yet implemented on host.
#endif
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_mipmap has been deprecated. "
                  "Instead use sample_mipmap.")
DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]],
                  const float level [[maybe_unused]]) {
  return sample_mipmap(imageHandle, coords, level);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_mipmap has been deprecated. "
                  "Instead use sample_mipmap.")
DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]],
                  const CoordT &dX [[maybe_unused]],
                  const CoordT &dY [[maybe_unused]]) {
  return sample_mipmap(imageHandle, coords, dX, dY);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                    const CoordT &coords [[maybe_unused]],
                    const float level [[maybe_unused]]) {
  detail::assert_sampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
                "Expected input coordinate to be have 1, 2, or 3 components "
                "for 1D, 2D and 3D images, respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to read a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(
        __invoke__ImageReadLod<HintT>(imageHandle.raw_handle, coords, level));
  }
#else
  assert(false); // Bindless images not yet implemented on host
#endif
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                    const CoordT &coords [[maybe_unused]],
                    const CoordT &dX [[maybe_unused]],
                    const CoordT &dY [[maybe_unused]]) {
  detail::assert_sampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
                "Expected input coordinates and gradients to have 1, 2, or 3 "
                "components for 1D, 2D, and 3D images, respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX,
                                          dY);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to read a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(
        __invoke__ImageReadGrad<HintT>(imageHandle.raw_handle, coords, dX, dY));
  }
#else
  assert(false); // Bindless images not yet implemented on host
#endif
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_image for mipmaps is deprecated. "
                  "Instead use sample_mipmap.")
DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 const CoordT &coords [[maybe_unused]],
                 const float level [[maybe_unused]]) {
  return sample_mipmap(imageHandle, coords, level);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
__SYCL_DEPRECATED("read_image for mipmaps is deprecated. "
                  "Instead use sample_mipmap.")
DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 const CoordT &coords [[maybe_unused]],
                 const CoordT &dX [[maybe_unused]],
                 const CoordT &dY [[maybe_unused]]) {
  return sample_mipmap(imageHandle, coords, dX, dY);
}
 
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT fetch_image_array(const unsampled_image_handle &imageHandle
                        [[maybe_unused]],
                        const CoordT &coords [[maybe_unused]],
                        const int arrayLayer [[maybe_unused]]) {
  detail::assert_unsampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2,
                "Expected input coordinate to be have 1 or 2 components for 1D "
                "and 2D images respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageArrayFetch<DataT>(imageHandle.raw_handle, coords,
                                            arrayLayer);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to fetch a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(__invoke__ImageArrayFetch<HintT>(
        imageHandle.raw_handle, coords, arrayLayer));
  }
#else
  assert(false); // Bindless images not yet implemented on host.
#endif
}
 
template <typename DataT>
DataT fetch_cubemap(const unsampled_image_handle &imageHandle,
                    const int2 &coords, const unsigned int face) {
  return fetch_image_array<DataT>(imageHandle, coords, face);
}
 
template <typename DataT, typename HintT = DataT>
DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]],
                     const sycl::float3 &dirVec [[maybe_unused]]) {
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    return __invoke__ImageReadCubemap<DataT, uint64_t>(imageHandle.raw_handle,
                                                       dirVec);
  } else {
    static_assert(sizeof(HintT) == sizeof(DataT),
                  "When trying to read a user-defined type, HintT must be of "
                  "the same size as the user-defined DataT.");
    static_assert(detail::is_recognized_standard_type<HintT>(),
                  "HintT must always be a recognized standard type");
    return sycl::bit_cast<DataT>(__invoke__ImageReadCubemap<DataT, uint64_t>(
        imageHandle.raw_handle, dirVec));
  }
#else
  assert(false); // Bindless images not yet implemented on host
#endif
}
 
template <typename DataT, typename CoordT>
void write_image(unsampled_image_handle imageHandle [[maybe_unused]],
                 const CoordT &coords [[maybe_unused]],
                 const DataT &color [[maybe_unused]]) {
  detail::assert_unsampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2 || coordSize == 3,
                "Expected input coordinate to be have 1, 2, or 3 components "
                "for 1D, 2D and 3D images, respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
  } else {
    // Convert DataT to a supported backend write type when user-defined type is
    // passed
    __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
                         detail::convert_color(color));
  }
#else
  assert(false); // Bindless images not yet implemented on host
#endif
}
 
template <typename DataT, typename CoordT>
void write_image_array(unsampled_image_handle imageHandle [[maybe_unused]],
                       const CoordT &coords [[maybe_unused]],
                       const int arrayLayer [[maybe_unused]],
                       const DataT &color [[maybe_unused]]) {
  detail::assert_unsampled_coords<CoordT>();
  constexpr size_t coordSize = detail::coord_size<CoordT>();
  static_assert(coordSize == 1 || coordSize == 2,
                "Expected input coordinate to be have 1 or 2 components for 1D "
                "and 2D images respectively.");
 
#ifdef __SYCL_DEVICE_ONLY__
  if constexpr (detail::is_recognized_standard_type<DataT>()) {
    __invoke__ImageArrayWrite(static_cast<uint64_t>(imageHandle.raw_handle),
                              coords, arrayLayer, color);
  } else {
    // Convert DataT to a supported backend write type when user-defined type is
    // passed
    __invoke__ImageArrayWrite(static_cast<uint64_t>(imageHandle.raw_handle),
                              coords, arrayLayer, detail::convert_color(color));
  }
#else
  assert(false); // Bindless images not yet implemented on host.
#endif
}
 
template <typename DataT>
void write_cubemap(unsampled_image_handle imageHandle, const sycl::int2 &coords,
                   const int face, const DataT &color) {
  return write_image_array(imageHandle, coords, face, color);
}
 
} // namespace ext::oneapi::experimental
} // namespace _V1
} // namespace sycl