[SYCL][Bindless][Exp] Allow read_image/mipmap to return user types (#11961)

- `read_image` and `read_mipmap` now allow user-defined types to be
returned.

- The `write_image` test case for user-defined types is updated to
include both reading and writing of user-defined types.

- An additional test is added for reading user-defined types where the
image type is a mipmap.

 - The extension document has been updated to reflect the above changes.

- The `write_image` function now checks whether a user-defined type was
passed to it, and converts the input data accordingly. If the type
passed has a specific NVPTX backend intrinsic, it will use that,
otherwise it will write the user data type bit-cast to a variant of
unsigned integer.

- This patch marks the bump to revision `5.0` of the bindless images
specification document.

Author: przemektmalon

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -962,10 +962,10 @@ listed above caused the failure.
 ```cpp
 namespace sycl::ext::oneapi::experimental {
 
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_image(const unsampled_image_handle &ImageHandle,
                  const CoordT &Coords);
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_image(const sampled_image_handle &ImageHandle, 
                  const CoordT &Coords);
 
@@ -981,26 +981,28 @@ will be fetched exactly as is in device memory. For the form that takes a
 `sampled_image_handle`, the image will be sampled according to the 
 `bindless_image_sampler` that was passed to the image upon construction.
 
-The returned data will be of templated type `DataT`, which is specified by the 
-user, and should map to the type that the image was created with (a combination 
-of `image_channel_type` and `image_channel_order`). One exception to this are 
-normalized integer channel types which are read back as either 32-bit or 16-bit 
-floating point values.
-
-For multi-channel types, the resultant `DataT` should be a `sycl::vec` type. 
-E.g., for a channel order of `image_channel_order::rg` and channel type of 
-`image_channel_type::fp16`, the resultant `DataT` should be 
-`sycl::vec<sycl::half, 2>`.
-
-An example of reading a normalized channel type with a channel order of 
-`image_channel_order::rg` and channel type of `image_channel_type::unorm_int8`, 
-the user can read the data as a `sycl::vec<float, 2>` or 
-`sycl::vec<sycl::half, 2>`.
+The user is required to pass a `DataT` template parameter, which specifies the
+return type of the `read_image` function. If `DataT` is not a recognized 
+standard type, as defined in <<recognized_standard_types>>, and instead a 
+user-defined type, the user must provide a `HintT` template parameter to the 
+`read_image` function, to allow the backend to select the correct device 
+intrinsic to fetch or sample their data.
+`HintT` must be one of the the <<recognized_standard_types>>, and must be the 
+same size as `DataT`.
+If `DataT` is a recognized standard type, and `HintT` is also passed, `HintT` 
+will be ignored.
+
+When reading a texture backed by a normalized integer channel type, either 
+`DataT` must be a 32-bit or 16-bit floating point value, a `sycl::vec` of 
+32-bit or 16-bit floating point values, or, in the case `DataT` is not one of 
+the above, then `HintT` must be one of the above, and be of the same size as 
+`DataT`.
 
 It's possible to write to an unsampled image via `write_image` passing the 
 handle of the image to be written to, along with the coordinates to write to and 
 the data. User-defined types are allowed to be written provided that type is 
-trivially copyable.
+trivially copyable. The user defined type must also be of the same size as any 
+of the <<recognized_standard_types>>.
 
 Sampled images cannot be written to using `write_image`.
 
@@ -1029,6 +1031,35 @@ Attempting to read an image with `read_mipmap` or any other defined read
 function will result in undefined behaviour.
 ====
 
+=== Recognized standard types [[recognized_standard_types]]
+
+For the purposes of this extension, the following are classified as recognized 
+standard types.
+
+* All POD types (`char`, `short`, `int`, `float`, etc.) excluding `double`
+* `sycl::half`
+* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `4`
+
+Any other types are classified as user-defined types.
+
+==== User-defined types
+
+Some examples of a user-defined types may be:
+
+```c++
+struct my_float4 {
+  float r, g, b, a;
+};
+
+struct my_short2 {
+  short r, g;
+};
+```
+
+When providing the above types as `DataT` parameters to an image read function, 
+the corresponding `HintT` parameters to use would be `sycl::vec<float, 4>` and 
+`sycl::vec<short, 2>`, respectively.
+
 == Mipmapped images
 
 So far, we have described how to create and operate on standard bindless images.
@@ -1128,13 +1159,13 @@ mipmap.
 
 ```c++
 // Nearest/linear filtering between mip levels
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_mipmap(const sampled_image_handle &ImageHandle,
                   const CoordT &Coords,
                   const float Level);
 
 // Anisotropic filtering
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_mipmap(const sampled_image_handle &ImageHandle,
                   const CoordT &Coords,
                   const CoordT &Dx, const CoordT &Dy);
@@ -1144,6 +1175,9 @@ Reading a mipmap follows the same restrictions on what coordinate types may be
 used as laid out in <<reading_writing_inside_kernel>>, and the viewing gradients 
 are bound to the same type as used for the coordinates.
 
+Reading a mipmap by providing a user-defined return `DataT` type also follows 
+the restrictions as laid out in <<reading_writing_inside_kernel>>.
+
 [NOTE]
 ====
 Attempting to read a mipmap with `read_image` or any other defined read function 
@@ -2017,4 +2051,8 @@ These features still need to be handled:
                     whether a `raw_sampler_handle` member is necessary.
                   - Renamed `image_handle` members in `sampled_image_handle` and
                     `unsampled_image_handle` structs to `raw_handle`.
+|5.0|2023-11-21| - Added section "Recognized standard types", to simplify 
+                   wording around what types are allowed to be read or written.
+                 - Allow `read_image` and `read_mipmap` to return a 
+                   user-defined type.
 |======================

sycl/include/sycl/ext/oneapi/bindless_images.hpp

@@ -742,12 +742,28 @@ template <typename CoordT> constexpr void assert_sampled_coords() {
                   "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
 
 /**
  *  @brief   Read an unsampled image using its handle
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. int, int2, or int4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The image handle
@@ -760,7 +776,7 @@ template <typename CoordT> constexpr void assert_sampled_coords() {
  *             The name mangling should therefore not interfere with one
  *             another
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]]) {
   detail::assert_unsampled_coords<CoordT>();
@@ -770,7 +786,17 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
+  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
@@ -780,6 +806,10 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
  *  @brief   Read a sampled image using its handle
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The image handle
@@ -792,7 +822,7 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
  *             The name mangling should therefore not interfere with one
  *             another
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                  const CoordT &coords [[maybe_unused]]) {
   detail::assert_sampled_coords<CoordT>();
@@ -802,7 +832,17 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
+  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
@@ -812,14 +852,18 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *  @brief   Read a mipmap image using its handle with LOD filtering
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The mipmap image handle
  *  @param   coords The coordinates at which to fetch mipmap image data
  *  @param   level The mipmap level at which to sample
  *  @return  Mipmap image data with LOD filtering
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                   const CoordT &coords [[maybe_unused]],
                   const float level [[maybe_unused]]) {
@@ -830,7 +874,17 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
+  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
@@ -840,6 +894,10 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
  *  @brief   Read a mipmap image using its handle with anisotropic filtering
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The mipmap image handle
@@ -848,7 +906,7 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
  *  @param   dY Screen space gradient in the y dimension
  *  @return  Mipmap image data with anisotropic filtering
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                   const CoordT &coords [[maybe_unused]],
                   const CoordT &dX [[maybe_unused]],
@@ -860,7 +918,18 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX, dY);
+  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
@@ -871,14 +940,18 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
  *           filtering
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The mipmap image handle
  *  @param   coords The coordinates at which to fetch mipmap image data
  *  @param   level The mipmap level at which to sample
  *  @return  Mipmap image data with LOD filtering
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 __SYCL_DEPRECATED("read_image for mipmaps is deprecated. "
                   "Instead use read_mipmap.")
 DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
@@ -891,7 +964,17 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
+  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
@@ -902,6 +985,10 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *           filtering
  *
  *  @tparam  DataT The return type
+ *  @tparam  HintT A hint type that can be used to select for a specialized
+ *           backend intrinsic when a user-defined type is passed as `DataT`.
+ *           HintT should be a `sycl::vec` type, `sycl::half` type, or POD type.
+ *           HintT must also have the same size as DataT.
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
  *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The mipmap image handle
@@ -910,7 +997,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *  @param   dY Screen space gradient in the y dimension
  *  @return  Mipmap image data with anisotropic filtering
  */
-template <typename DataT, typename CoordT>
+template <typename DataT, typename HintT = DataT, typename CoordT>
 __SYCL_DEPRECATED("read_image for mipmaps is deprecated. "
                   "Instead use read_mipmap.")
 DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
@@ -924,7 +1011,18 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-  return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX, dY);
+  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
@@ -951,8 +1049,14 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
-  __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
-                       detail::convert_color_nvptx(color));
+  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_nvptx(color));
+  }
 #else
   __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
 #endif