[SYCL][Bindless][E2E] Enable 3-channel image test for Intel GPUs

[przemektmalon]

This patch enables the 3-channel image E2E test on Intel GPUs.

[wenju-he]

@przemektmalon could you please update test to use ushort format like in following diff

diff --git a/sycl/test-e2e/bindless_images/3_channel_format.cpp b/sycl/test-e2e/bindless_images/3_channel_format.cpp
index a3668f4f3197..4a081c95ed4a 100644
--- a/sycl/test-e2e/bindless_images/3_channel_format.cpp
+++ b/sycl/test-e2e/bindless_images/3_channel_format.cpp
@@ -21,19 +21,19 @@ int main() {
   auto ctxt = q.get_context();

   constexpr size_t width = 512;
-  std::vector<float> out(width);
-  std::vector<float> expected(width);
-  std::vector<sycl::float3> dataIn(width);
-  float exp = 512;
-  for (int i = 0; i < width; i++) {
+  std::vector<unsigned short> out(width);
+  std::vector<unsigned short> expected(width);
+  std::vector<sycl::ushort3> dataIn(width);
+  unsigned short exp = 512;
+  for (unsigned i = 0; i < width; i++) {
     expected[i] = exp;
-    dataIn[i] = sycl::float3(exp, width, i);
+    dataIn[i] = sycl::ushort3(exp, width, i);
   }

   try {
     // Main point of this test is to check creating an image
     // with a 3-channel format
-    syclexp::image_descriptor desc({width}, 3, sycl::image_channel_type::fp32);
+    syclexp::image_descriptor desc({width}, 3, sycl::image_channel_type::unsigned_int16);

     syclexp::image_mem imgMem(desc, dev, ctxt);

@@ -46,7 +46,7 @@ int main() {
     syclexp::unsampled_image_handle imgHandle =
         sycl::ext::oneapi::experimental::create_image(imgMem, desc, dev, ctxt);

-    sycl::buffer<float> buf(out.data(), width);
+    sycl::buffer<unsigned short> buf(out.data(), width);

     q.submit([&](sycl::handler &cgh) {
       sycl::accessor outAcc{buf, cgh};
@@ -55,9 +55,9 @@ int main() {
 #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__)
         // This shouldn't be hit anyway since CUDA doesn't support
         // 3-channel formats, but we need to ensure the kernel can compile
-        using pixel_t = sycl::float4;
+        using pixel_t = sycl::ushort4;
 #else
-        using pixel_t = sycl::float3;
+        using pixel_t = sycl::ushort3;
 #endif
         auto pixel = syclexp::fetch_image<pixel_t>(imgHandle, int(id[0]));
         outAcc[id] = pixel[0];

Intel GPU currently only supports ushort and uchar format for 3-channels.

[wenju-he]

std::vector<sycl::ushort3> dataIn(width); should be changed to std::vector<sycl::ushort> dataIn(width * 3); because sycl::ushort3 has the same size as sycl::ushort4 and intel gpu NEO runtime assumes memory layout of 3-channel image is of 3-component per element. Thanks @AshwinKumarKulkarni for pointing this out.

[przemektmalon]

std::vector<sycl::ushort3> dataIn(width); should be changed to std::vector<sycl::ushort> dataIn(width * 3); because sycl::ushort3 has the same size as sycl::ushort4 and intel gpu NEO runtime assumes memory layout of 3-channel image is of 3-component per element. Thanks @AshwinKumarKulkarni for pointing this out.

Indeed, good catch. Updated to use unsigned short, as sycl::ushort throws compilation errors due to deprecation.

@wenju-he Does this mean the kernel code requires change as well? I.e. when calling syclexp::fetch_image<sycl::ushort3>?

If the Intel driver returns a 6-byte length type, then we might need to use a custom ushort3 struct in the test which satisfies the size requirements, as well as updating the is_data_size_valid() function in bindless_images.hpp here.