[SYCL][Bindless] Enable SPIRV path for bindless_images (#11886)

This is 1st patch to enable SPIRV path for bindless_images so that we
can start
generating SPIRV for testing purpose.

Author: jsji

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

@@ -761,11 +761,7 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
-#else
-  // TODO: add SPIRV part for unsampled image read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host
 #endif
@@ -797,11 +793,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
-#else
-  // TODO: add SPIRV part for sampled image read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host.
 #endif
@@ -829,11 +821,7 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
-#else
-  // TODO: add SPIRV for mipmap level read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host
 #endif
@@ -863,11 +851,7 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX, dY);
-#else
-  // TODO: add SPIRV part for mipmap grad read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host
 #endif
@@ -898,11 +882,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "for 1D, 2D and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
-#else
-  // TODO: add SPIRV for mipmap level read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host
 #endif
@@ -935,11 +915,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]],
                 "components for 1D, 2D, and 3D images, respectively.");
 
 #ifdef __SYCL_DEVICE_ONLY__
-#if defined(__NVPTX__)
   return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX, dY);
-#else
-  // TODO: add SPIRV part for mipmap grad read
-#endif
 #else
   assert(false); // Bindless images not yet implemented on host
 #endif
@@ -969,7 +945,7 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
   __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
                        detail::convert_color_nvptx(color));
 #else
-  // TODO: add SPIRV part for unsampled image write
+  __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
 #endif
 #else
   assert(false); // Bindless images not yet implemented on host

sycl/test/extensions/bindless_images.cpp

@@ -0,0 +1,51 @@
+// RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s
+
+#include <iostream>
+#include <sycl/sycl.hpp>
+
+// CHECK: spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI10image_readEE
+// CHECK: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4
+using namespace sycl::ext::oneapi::experimental;
+class image_read;
+int main() {
+
+  sycl::device dev;
+  sycl::queue q(dev);
+  auto ctxt = q.get_context();
+
+  constexpr size_t width = 512;
+  std::vector<float> out(width);
+  std::vector<sycl::float4> dataIn1(width);
+  for (int i = 0; i < width; i++) {
+    dataIn1[i] = sycl::float4(i, i, i, i);
+  }
+
+  {
+    image_descriptor desc({width}, sycl::image_channel_order::rgba,
+                          sycl::image_channel_type::fp32);
+
+    image_mem imgMem0(desc, dev, ctxt);
+    unsampled_image_handle imgHandle1 = create_image(imgMem0, desc, dev, ctxt);
+
+    q.ext_oneapi_copy(dataIn1.data(), imgMem0.get_handle(), desc);
+    q.wait_and_throw();
+
+    sycl::buffer<float, 1> buf((float *)out.data(), width);
+    q.submit([&](sycl::handler &cgh) {
+      auto outAcc = buf.get_access<sycl::access_mode::write>(cgh, width);
+
+      cgh.parallel_for<image_read>(width, [=](sycl::id<1> id) {
+        sycl::float4 px1 = read_image<sycl::float4>(imgHandle1, int(id[0]));
+        outAcc[id] = px1[0];
+      });
+    });
+
+    q.wait_and_throw();
+    destroy_image_handle(imgHandle1, dev, ctxt);
+  }
+
+  for (int i = 0; i < width; i++) {
+    std::cout << "Actual: " << out[i] << std::endl;
+  }
+  return 0;
+}