[Bindless][SYCL][Doc] Add HintT tparam to cubemap fetch and sample (#13742)

Add HintT template parameter to the fetch_cubemap and sample_cubemap
funcs in the bindless spec.

Add HintT template parameter and doxygen to the fetch_cubemap and
sample_cubemap implementation.

Modify existing cubemap_sampled.cpp test to test the sampling with a
user-defined type.

Author: Seanst98

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -1506,13 +1506,13 @@ sampling depends on the sampler attributes passed upon creation of the cubemap.
 
 ```c++
 // Unsampled cubemap read 
-template <typename DataT>
+template <typename DataT, typename HintT = DataT>
 DataT fetch_cubemap(const unsampled_image_handle &ImageHandle,
                     const int2 &Coords,
                     const int Face);
 
 // Sampled cubemap read
-template <typename DataT>
+template <typename DataT, typename HintT = DataT>
 DataT sample_cubemap(const sampled_image_handle &ImageHandle,
                      const float3 &Vec);
 
@@ -2684,4 +2684,6 @@ These features still need to be handled:
                    parameter. 
 |5.7|2024-04-09| - Allow fetching of sampled image data through the 
                    `fetch_image` API.
+|5.8|2024-05-09| - Add missing cubemap `HintT` template parameter to 
+                   `fetch_cubemap` and `sample_cubemap`.
 |======================

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

@@ -1244,22 +1244,30 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle
  *  @brief   Fetch data from an unsampled cubemap 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.
  *
  *  @param   imageHandle The image handle
  *  @param   coords The coordinates at which to fetch image data (int2 only)
  *  @param   face The cubemap face at which to fetch
  *  @return  Image data
  */
-template <typename DataT>
+template <typename DataT, typename HintT = DataT>
 DataT fetch_cubemap(const unsampled_image_handle &imageHandle,
                     const int2 &coords, const unsigned int face) {
-  return fetch_image_array<DataT>(imageHandle, coords, face);
+  return fetch_image_array<DataT, HintT>(imageHandle, coords, face);
 }
 
 /**
  *  @brief   Sample a cubemap 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.
  *
  *  @param   imageHandle The image handle
  *  @param   dirVec The direction vector at which to sample image data (float3
@@ -1280,7 +1288,7 @@ DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]],
                   "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>(
+    return sycl::bit_cast<DataT>(__invoke__ImageReadCubemap<HintT, uint64_t>(
         imageHandle.raw_handle, dirVec));
   }
 #else

sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp

@@ -4,6 +4,7 @@
 // RUN: %{build} -o %t.out
 // RUN: %{run} %t.out
 
+#include "../user_types/user_types_common.hpp"
 #include <iostream>
 #include <sycl/sycl.hpp>
 
@@ -27,12 +28,12 @@ int main() {
   size_t N = width * height;
   std::vector<float> out(N);
   std::vector<float> expected(N);
-  std::vector<sycl::float4> dataIn1(N * 6);
+  std::vector<sycl::float4> dataIn(N * 6);
   for (int i = 0; i < width; i++) {
     for (int j = 0; j < height; j++) {
       for (int k = 0; k < 6; k++) {
-        dataIn1[i + width * (j + height * k)] = {i + width * (j + height * k),
-                                                 0, 0, 0};
+        dataIn[i + width * (j + height * k)] = {i + width * (j + height * k), 0,
+                                                0, 0};
       }
     }
   }
@@ -63,7 +64,7 @@ int main() {
 
     // Extension: copy over data to device (handler variant).
     q.submit([&](sycl::handler &cgh) {
-      cgh.ext_oneapi_copy(dataIn1.data(), imgMem.get_handle(), desc);
+      cgh.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc);
     });
     q.wait_and_throw();
 
@@ -95,11 +96,13 @@ int main() {
             float fdim2 = (((float(dim1) / (float)height) * 1.98) - 0.99) +
                           (1.f / (float)height);
 
+            my_float4 myPixel{};
+
             // Extension: sample cubemap data from handle.
-            sycl::float4 px = syclexp::sample_cubemap<sycl::float4>(
+            myPixel = syclexp::sample_cubemap<my_float4, sycl::float4>(
                 imgHandle, sycl::float3(fdim0, fdim1, fdim2));
 
-            outAcc[sycl::id<2>{dim0, dim1}] = px[0];
+            outAcc[sycl::id<2>{dim0, dim1}] = myPixel.x;
           });
     });
     q.wait_and_throw();