[SYCL] Fix sampled bindless fetch_image for float type (#20107)

The implementation previously used `__spirv_ImageSampleExplicitLod`
which returned incorrect results when the `DataT` return type was a
vector of floats. This change makes it use `__spirv_ImageRead` instead
and adds a test case for the `float` image type.

Signed-off-by: Michael Aziz <[email protected]>

Author: 0x12CC

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

@@ -1036,15 +1036,19 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "HintT must always be a recognized standard type");
 
 #ifdef __SYCL_DEVICE_ONLY__
+  // Convert the raw handle to an image and use FETCH_UNSAMPLED_IMAGE since
+  // fetch_image should not use the sampler
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return FETCH_SAMPLED_IMAGE(
+    return FETCH_UNSAMPLED_IMAGE(
         DataT,
-        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
         coords);
   } else {
-    return sycl::bit_cast<DataT>(FETCH_SAMPLED_IMAGE(
+    return sycl::bit_cast<DataT>(FETCH_UNSAMPLED_IMAGE(
         HintT,
-        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
         coords));
   }
 #else

sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp

@@ -11,9 +11,10 @@
 #include <sycl/ext/oneapi/bindless_images.hpp>
 #include <sycl/usm.hpp>
 
-class kernel_sampled_fetch;
+namespace {
 
-int main() {
+template <typename T, sycl::image_channel_type ChanType>
+static int testSampledImageFetch() {
 
   sycl::device dev;
   sycl::queue q(dev);
@@ -23,9 +24,9 @@ int main() {
   constexpr size_t width = 5;
   constexpr size_t height = 6;
   constexpr size_t N = width * height;
-  std::vector<sycl::vec<uint16_t, 4>> out(N);
-  std::vector<sycl::vec<uint16_t, 4>> expected(N);
-  std::vector<sycl::vec<uint16_t, 4>> dataIn(N);
+  std::vector<sycl::vec<T, 4>> out(N);
+  std::vector<sycl::vec<T, 4>> expected(N);
+  std::vector<sycl::vec<T, 4>> dataIn(N);
   for (int i = 0; i < width; i++) {
     for (int j = 0; j < height; j++) {
       auto index = i + (width * j);
@@ -43,8 +44,7 @@ int main() {
         sycl::filtering_mode::linear);
 
     // Extension: image descriptor
-    syclexp::image_descriptor desc({width, height}, 4,
-                                   sycl::image_channel_type::unsigned_int16);
+    syclexp::image_descriptor desc({width, height}, 4, ChanType);
     size_t pitch = 0;
 
     // Extension: returns the device pointer to USM allocated pitched memory
@@ -65,21 +65,20 @@ int main() {
 
     sycl::buffer buf(out.data(), sycl::range{height, width});
     q.submit([&](sycl::handler &cgh) {
-      auto outAcc = buf.get_access<sycl::access_mode::write>(
+      auto outAcc = buf.template get_access<sycl::access_mode::write>(
           cgh, sycl::range<2>{height, width});
 
-      cgh.parallel_for<kernel_sampled_fetch>(
-          sycl::nd_range<2>{{width, height}, {width, height}},
-          [=](sycl::nd_item<2> it) {
-            size_t dim0 = it.get_local_id(0);
-            size_t dim1 = it.get_local_id(1);
+      cgh.parallel_for(sycl::nd_range<2>{{width, height}, {width, height}},
+                       [=](sycl::nd_item<2> it) {
+                         size_t dim0 = it.get_local_id(0);
+                         size_t dim1 = it.get_local_id(1);
 
-            // Extension: fetch data from sampled image handle
-            auto px1 = syclexp::fetch_image<sycl::vec<uint16_t, 4>>(
-                imgHandle, sycl::int2(dim0, dim1));
+                         // Extension: fetch data from sampled image handle
+                         auto px1 = syclexp::fetch_image<sycl::vec<T, 4>>(
+                             imgHandle, sycl::int2(dim0, dim1));
 
-            outAcc[sycl::id<2>{dim1, dim0}] = px1;
-          });
+                         outAcc[sycl::id<2>{dim1, dim0}] = px1;
+                       });
     });
 
     q.wait_and_throw();
@@ -121,3 +120,23 @@ int main() {
   std::cout << "Test failed!" << std::endl;
   return 3;
 }
+
+} // namespace
+
+int main() {
+  if (int err =
+          testSampledImageFetch<uint16_t,
+                                sycl::image_channel_type::unsigned_int16>()) {
+    return err;
+  }
+  if (int err =
+          testSampledImageFetch<uint32_t,
+                                sycl::image_channel_type::unsigned_int32>()) {
+    return err;
+  }
+  if (int err =
+          testSampledImageFetch<float, sycl::image_channel_type::fp32>()) {
+    return err;
+  }
+  return 0;
+}