[SYCL][E2E][Bindless] Add 3-channel unorm_int8 and half vulkan-interop subtests for L0 (intel#16374)

Also disable failing unsampled unorm tests on intel gpu.

Author: wenju-he

sycl/test-e2e/bindless_images/dx12_interop/read_write_unsampled.cpp

@@ -3,7 +3,7 @@
 // REQUIRES: build-and-run-mode
 
 // DEFINE: %{link-flags}=%if cl_options %{ /clang:-ld3d12 /clang:-ldxgi /clang:-ldxguid %} %else %{ -ld3d12 -ldxgi -ldxguid %}
-// RUN: %{build} %{link-flags} -o %t.out
+// RUN: %{build} %{link-flags} -o %t.out %if any-device-is-level_zero %{ -DDISABLE_UNORM_TESTS %}
 // RUN: %{run-unfiltered-devices} env NEOReadDebugKeys=1 UseBindlessMode=1 UseExternalAllocatorForSshAndDsh=1 %t.out
 
 #pragma clang diagnostic ignored "-Waddress-of-temporary"
@@ -733,8 +733,10 @@ int main() {
   validated &=
       runTest<1, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
                               globalSize1, localSize1);
+#ifndef DISABLE_UNORM_TESTS
   validated &= runTest<1, uint8_t, 4>(
       device, sycl::image_channel_type::unorm_int8, globalSize1, localSize1);
+#endif
   validated &= runTest<1, float, 1>(device, sycl::image_channel_type::fp32,
                                     globalSize1, localSize1);
   validated &= runTest<1, sycl::half, 2>(device, sycl::image_channel_type::fp16,
@@ -752,8 +754,10 @@ int main() {
   validated &=
       runTest<2, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
                               globalSize2[0], {16, 16});
+#ifndef DISABLE_UNORM_TESTS
   validated &= runTest<2, uint8_t, 4>(
       device, sycl::image_channel_type::unorm_int8, globalSize2[1], {16, 8});
+#endif
   validated &= runTest<2, float, 1>(device, sycl::image_channel_type::fp32,
                                     globalSize2[2], {16, 8});
   validated &= runTest<2, sycl::half, 2>(device, sycl::image_channel_type::fp16,
@@ -774,8 +778,10 @@ int main() {
   validated &=
       runTest<3, uint32_t, 1>(device, sycl::image_channel_type::unsigned_int32,
                               globalSize3[0], {16, 16, 1});
+#ifndef DISABLE_UNORM_TESTS
   validated &= runTest<3, uint8_t, 4>(
       device, sycl::image_channel_type::unorm_int8, globalSize3[1], {16, 8, 2});
+#endif
   validated &= runTest<3, float, 1>(device, sycl::image_channel_type::fp32,
                                     globalSize3[2], {16, 8, 1});
   validated &= runTest<3, sycl::half, 2>(device, sycl::image_channel_type::fp16,

sycl/test-e2e/bindless_images/helpers/common.hpp

@@ -82,6 +82,8 @@ constexpr sycl::vec<DType, NChannel> init_vector(DType val) {
     return sycl::vec<DType, NChannel>{val};
   } else if constexpr (NChannel == 2) {
     return sycl::vec<DType, NChannel>{val, val};
+  } else if constexpr (NChannel == 3) {
+    return sycl::vec<DType, NChannel>{val, val, val};
   } else if constexpr (NChannel == 4) {
     return sycl::vec<DType, NChannel>{val, val, val, val};
   } else {

sycl/test-e2e/bindless_images/vulkan_interop/sampled_images.cpp

@@ -483,6 +483,12 @@ bool run_tests() {
   valid &= run_test<2, sycl::half, 2, sycl::image_channel_type::fp16,
                     sycl::image_channel_order::rg, class fp16_2d_c2>(
       {1920, 1080}, {16, 8}, 0);
+  valid &= run_test<2, sycl::half, 3, sycl::image_channel_type::fp16,
+                    sycl::image_channel_order::rgb, class fp16_2d_c3>(
+      {2048, 2048}, {16, 16}, 0);
+  valid &= run_test<2, uint8_t, 3, sycl::image_channel_type::unorm_int8,
+                    sycl::image_channel_order::rgb, class unorm_int8_2d_c3>(
+      {2048, 2048}, {16, 16}, 0);
   valid &= run_test<2, sycl::half, 4, sycl::image_channel_type::fp16,
                     sycl::image_channel_order::rgba, class fp16_2d_c4>(
       {2048, 2048}, {16, 16}, 0);

sycl/test-e2e/bindless_images/vulkan_interop/unsampled_images.cpp

@@ -590,11 +590,6 @@ bool run_all() {
                     sycl::image_channel_order::rgba, class fp16_3d_c4>(
       {2048, 2048, 4}, {16, 16, 1}, seed);
 
-  printString("Running 3D unorm_int8_c4\n");
-  valid &= run_test<3, uint8_t, 4, sycl::image_channel_type::unorm_int8,
-                    sycl::image_channel_order::rgba, class unorm_int8_3d_c4>(
-      {2048, 2048, 2}, {16, 16, 1}, seed);
-
   printString("Running 2D float\n");
   valid &= run_test<2, float, 1, sycl::image_channel_type::fp32,
                     sycl::image_channel_order::r, class fp32_2d_c1>(
@@ -608,10 +603,12 @@ bool run_all() {
                     sycl::image_channel_order::rgba, class fp16_2d_c4>(
       {2048, 2048}, {16, 16}, seed);
 
-  printString("Running 2D unorm_int8_c4\n");
-  valid &= run_test<2, uint8_t, 4, sycl::image_channel_type::unorm_int8,
-                    sycl::image_channel_order::rgba, class unorm_int8_2d_c4>(
+  // 3-channels
+  printString("Running 2D half3\n");
+  valid &= run_test<2, sycl::half, 3, sycl::image_channel_type::fp16,
+                    sycl::image_channel_order::rgb, class fp16_2d_c3>(
       {2048, 2048}, {2, 2}, seed);
+
 #else
   printString("Running 3D uint4\n");
   valid &= run_test<3, uint32_t, 4, sycl::image_channel_type::signed_int32,

sycl/test-e2e/bindless_images/vulkan_interop/vulkan_common.hpp

@@ -827,6 +827,8 @@ VkFormat to_vulkan_format(sycl::image_channel_order order,
       return VK_FORMAT_R8_UNORM;
     case sycl::image_channel_order::rg:
       return VK_FORMAT_R8G8_UNORM;
+    case sycl::image_channel_order::rgb:
+      return VK_FORMAT_R8G8B8_UNORM;
     case sycl::image_channel_order::rgba:
       return VK_FORMAT_R8G8B8A8_UNORM;
     default: {
@@ -894,6 +896,8 @@ VkFormat to_vulkan_format(sycl::image_channel_order order,
       return VK_FORMAT_R16_SFLOAT;
     case sycl::image_channel_order::rg:
       return VK_FORMAT_R16G16_SFLOAT;
+    case sycl::image_channel_order::rgb:
+      return VK_FORMAT_R16G16B16_SFLOAT;
     case sycl::image_channel_order::rgba:
       return VK_FORMAT_R16G16B16A16_SFLOAT;
     default: {