[Bindless][SYCL] read/write images with only acceptable coordinate types (#11341)

Assert the restrictions on reading/writing coordinate types and mipmap
anisotropic viewing gradient types

Author: Seanst98

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -945,7 +945,7 @@ with error code `sycl::errc::invalid`, and relay an error message back to the
 user through `sycl::exception::what()`, describing which of the scenarios 
 listed above caused the failure.
 
-=== Reading and writing inside the kernel
+=== Reading and writing inside the kernel [[reading_writing_inside_kernel]]
 
 ```cpp
 namespace sycl::ext::oneapi::experimental {
@@ -994,16 +994,16 @@ trivially copyable.
 
 Sampled images cannot be written to using `write_image`.
 
-For unsampled images, coordinates are specified by `int`, `sycl::vec<int, 2>`, 
-and `sycl::vec<int, 4>` for 1D, 2D, and 3D images respectively.
+For reading and writing of unsampled images, coordinates are specified by `int`, 
+`sycl::vec<int, 2>`, and `sycl::vec<int, 4>` for 1D, 2D, and 3D images, 
+respectively.
 
 Sampled image reads take `float`, `sycl::vec<float, 2>`, and 
-`sycl::vec<float, 4>` coordinate types for 1D, 2D, and 3D images respectively. 
+`sycl::vec<float, 4>` coordinate types for 1D, 2D, and 3D images, respectively.
 
-In the case of 3D reads or writes, the fourth element in the coordinate vector 
-is ignored.
-
-Note that coordinates for 3D images take a vector of size 4, not 3.
+Note that in the case of 3D reads or writes, coordinates for 3D images take a 
+vector of size 4, not 3, as the fourth element in the coordinate vector is 
+ignored.
 
 Note also that all images must be used in either read-only or write-only fashion 
 within a single kernel invocation; read/write images are not supported.
@@ -1124,6 +1124,10 @@ DataT read_image(const sampled_image_handle &ImageHandle,
                  const CoordT &Dx, const CoordT &Dy);
 ```
 
+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.
+
 == Interoperability
 
 === Querying interoperability support
@@ -1976,4 +1980,5 @@ These features still need to be handled:
                    (`unorm_short_555`, `unorm_short_565`, `unorm_int_101010`)
 |4.4|2023-09-12| - Added overload with `sycl::queue` to standalone functions
 |4.5|2023-09-14| - Update wording for allocating images + fix typo
+|4.6|2023-09-19| - Clarify restrictions on reading/writing coordinate types
 |======================

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

@@ -633,15 +633,26 @@ 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<CoordT>::value) {
+  if constexpr (std::is_scalar_v<CoordT>) {
     return 1;
   } else {
     return CoordT::size();
   }
 }
 
+#if defined(__NVPTX__)
 // bit_cast Color to a type the NVPTX backend is known to accept
 template <typename DataT> constexpr auto convert_color_nvptx(DataT Color) {
   constexpr size_t dataSize = sizeof(DataT);
@@ -662,15 +673,39 @@ template <typename DataT> constexpr auto convert_color_nvptx(DataT Color) {
     return sycl::bit_cast<sycl::vec<uint32_t, 4>>(Color);
   }
 }
+#endif
+
+// 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");
+  }
+}
 } // namespace detail
 
 /**
  *  @brief   Read an unsampled image using its handle
  *
  *  @tparam  DataT The return type
  *  @tparam  CoordT The input coordinate type. e.g. int, int2, or int4 for
- *           1D, 2D, and 3D respectively
+ *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The image handle
  *  @param   coords The coordinates at which to fetch image data
  *  @return  Image data
@@ -684,10 +719,11 @@ template <typename DataT> constexpr auto convert_color_nvptx(DataT Color) {
 template <typename DataT, typename CoordT>
 DataT read_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 == 4,
                 "Expected input coordinate to be have 1, 2, or 4 components "
-                "for 1D, 2D and 3D images respectively.");
+                "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
@@ -705,7 +741,7 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
  *
  *  @tparam  DataT The return type
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
- *           1D, 2D, and 3D respectively
+ *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The image handle
  *  @param   coords The coordinates at which to fetch image data
  *  @return  Sampled image data
@@ -719,10 +755,11 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
 template <typename DataT, typename CoordT>
 DataT read_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 == 4,
                 "Expected input coordinate to be have 1, 2, or 4 components "
-                "for 1D, 2D and 3D images respectively.");
+                "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
@@ -740,7 +777,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *
  *  @tparam  DataT The return type
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
- *           1D, 2D, and 3D respectively
+ *           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
@@ -750,10 +787,11 @@ template <typename DataT, typename CoordT>
 DataT read_image(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 == 4,
                 "Expected input coordinate to be have 1, 2, or 4 components "
-                "for 1D, 2D and 3D images respectively.");
+                "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
@@ -771,7 +809,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *
  *  @tparam  DataT The return type
  *  @tparam  CoordT The input coordinate type. e.g. float, float2, or float4 for
- *           1D, 2D, and 3D respectively
+ *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The mipmap image handle
  *  @param   coords The coordinates at which to fetch mipmap image data
  *  @param   dX Screen space gradient in the x dimension
@@ -783,11 +821,11 @@ 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]]) {
+  detail::assert_sampled_coords<CoordT>();
   constexpr size_t coordSize = detail::coord_size<CoordT>();
   static_assert(coordSize == 1 || coordSize == 2 || coordSize == 4,
-                "Expected input coordinate and gradient to have 1, 2, or 4 "
-                "components "
-                "for 1D, 2D and 3D images respectively.");
+                "Expected input coordinates and gradients to have 1, 2, or 4 "
+                "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
@@ -805,23 +843,24 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
  *
  *  @tparam  DataT The data type to write
  *  @tparam  CoordT The input coordinate type. e.g. int, int2, or int4 for
- *           1D, 2D, and 3D respectively
+ *           1D, 2D, and 3D, respectively
  *  @param   imageHandle The image handle
  *  @param   coords The coordinates at which to write image data
  */
 template <typename DataT, typename CoordT>
 void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
-                 const CoordT &Coords [[maybe_unused]],
-                 const DataT &Color [[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 == 4,
                 "Expected input coordinate to be have 1, 2, or 4 components "
-                "for 1D, 2D and 3D images respectively.");
+                "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
 #if defined(__NVPTX__)
-  __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, Coords,
-                       detail::convert_color_nvptx(Color));
+  __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
+                       detail::convert_color_nvptx(color));
 #else
   // TODO: add SPIRV part for unsampled image write
 #endif