[SYCL][Bindless] Update spirv read/fetch from sampled image and sampled image array (intel#14493)

For spirv path, call __invoke__ImageReadLod to generate image builtin
represented as ImageSampleExplicitLod instruction in SPIRV that can be
consumed by Intel Graphics Compiler.

Author: wenju-he

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

@@ -812,6 +812,11 @@ template <int NDims>
 using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type<
     NDims, sycl::access::mode::write, sycl::access::target::image_array>::type;
 
+template <int NDims>
+using OCLSampledImageArrayTyRead =
+    typename sycl::detail::sampled_opencl_image_type<
+        detail::OCLImageArrayTyRead<NDims>>::type;
+
 // Macros are required because it is not legal for a function to return
 // a variable of type 'opencl_image_type'.
 #if defined(__SPIR__)
@@ -823,34 +828,64 @@ using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type<
       typename sycl::detail::sampled_opencl_image_type<                        \
           detail::OCLImageTyRead<NDims>>::type>(raw_handle)
 
+#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims)               \
+  __spirv_ConvertHandleToSampledImageINTEL<                                    \
+      typename sycl::detail::sampled_opencl_image_type<                        \
+          detail::OCLImageArrayTyRead<NDims>>::type>(raw_handle)
+
 #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
   __invoke__ImageRead<DataT>(raw_handle, coords)
 
 #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords)                         \
-  __invoke__ImageRead<DataT>(raw_handle, coords)
+  __invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)
+
+#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords)                           \
+  __invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)
 
 #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer)  \
   __invoke__ImageRead<DataT>(raw_handle, coordsLayer)
 
 #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color)  \
   __invoke__ImageWrite(raw_handle, coordsLayer, color)
 
+#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer,       \
+                                  coordsLayer)                                 \
+  __invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)
+
+#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer,        \
+                                 coordsLayer)                                  \
+  __invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)
+
 #else
 #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle
 
 #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle
 
+#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) raw_handle
+
 #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
   __invoke__ImageFetch<DataT>(raw_handle, coords)
 
 #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords)                         \
   __invoke__SampledImageFetch<DataT>(raw_handle, coords)
 
+#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords)                           \
+  __invoke__ImageRead<DataT>(raw_handle, coords)
+
 #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer)  \
   __invoke__ImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)
 
 #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color)  \
   __invoke__ImageArrayWrite(raw_handle, coords, arrayLayer, color)
+
+#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer,       \
+                                  coordsLayer)                                 \
+  __invoke__SampledImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)
+
+#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer,        \
+                                 coordsLayer)                                  \
+  __invoke__ImageArrayRead<DataT>(raw_handle, coords, arrayLayer)
+
 #endif
 
 #endif // __SYCL_DEVICE_ONLY__
@@ -1030,11 +1065,13 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageRead<DataT>(
+    return SAMPLE_IMAGE_READ(
+        DataT,
         CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
         coords);
   } else {
-    return sycl::bit_cast<DataT>(__invoke__ImageRead<HintT>(
+    return sycl::bit_cast<DataT>(SAMPLE_IMAGE_READ(
+        HintT,
         CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
         coords));
   }
@@ -1410,17 +1447,23 @@ DataT fetch_image_array(const sampled_image_handle &imageHandle
                 "and 2D images respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
+  sycl::vec<int, coordSize + 1> coordsLayer{coords, arrayLayer};
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__SampledImageArrayFetch<DataT>(imageHandle.raw_handle,
-                                                   coords, arrayLayer);
+    return FETCH_SAMPLED_IMAGE_ARRAY(DataT,
+                                     CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
+                                         imageHandle.raw_handle, coordSize),
+                                     coords, arrayLayer, coordsLayer);
   } 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__SampledImageArrayFetch<HintT>(
-        imageHandle.raw_handle, coords, arrayLayer));
+    return sycl::bit_cast<DataT>(
+        FETCH_SAMPLED_IMAGE_ARRAY(HintT,
+                                  CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
+                                      imageHandle.raw_handle, coordSize),
+                                  coords, arrayLayer, coordsLayer));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.
@@ -1454,17 +1497,23 @@ DataT sample_image_array(const sampled_image_handle &imageHandle
                 "and 2D images respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
+  sycl::vec<float, coordSize + 1> coordsLayer{coords, arrayLayer};
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageArrayRead<DataT>(imageHandle.raw_handle, coords,
-                                           arrayLayer);
+    return READ_SAMPLED_IMAGE_ARRAY(DataT,
+                                    CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
+                                        imageHandle.raw_handle, coordSize),
+                                    coords, arrayLayer, coordsLayer);
   } 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__ImageArrayRead<HintT>(
-        imageHandle.raw_handle, coords, arrayLayer));
+    return sycl::bit_cast<DataT>(
+        READ_SAMPLED_IMAGE_ARRAY(HintT,
+                                 CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
+                                     imageHandle.raw_handle, coordSize),
+                                 coords, arrayLayer, coordsLayer));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.

sycl/test/check_device_code/extensions/bindless_images.cpp

@@ -4,7 +4,7 @@
 // RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-LLVM
 
 // CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4
-// CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4
+// CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z30__spirv_ImageSampleExplicitLod
 // CHECK-LLVM: tail call spir_func void @_Z18__spirv_ImageWriteI14
 
 // RUN: %clangxx -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o %t.out
@@ -50,7 +50,7 @@
 
 // Read sampled image
 // Arguments: Result Type, Result, Image, Coords
-// CHECK-SPIRV-NEXT: ImageRead [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}}
+// CHECK-SPIRV-NEXT: ImageSampleExplicitLod [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}}
 
 // Convert handle to SPIR-V image
 // Arguments: Result Type, Result, Handle