[SYCL][E2E][Bindless] Fix access to mapped memory of 3-channel vulkan image (#16899)

We can't use sycl::vec<T, 3> to access mapped memory of 3-channel vulkan
image because the sycl type is aligned to 4 elements but there are only
3 elements per pixel in vulkan image memory.

Author: wenju-he

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

@@ -102,10 +102,12 @@ bool equal_vec(sycl::vec<DType, NChannels> v1, sycl::vec<DType, NChannels> v2) {
   return true;
 }
 
-template <typename DType, int NChannels>
-static void fill_rand(std::vector<sycl::vec<DType, NChannels>> &v,
+template <typename T>
+static void fill_rand(std::vector<T> &v,
                       int seed = std::default_random_engine::default_seed) {
   assert(!v.empty());
+  using DType = sycl::detail::get_elem_type_t<T>;
+  constexpr int NChannels = sycl::detail::get_vec_size<T>::size;
   std::default_random_engine generator;
   generator.seed(seed);
   auto distribution = [&]() {
@@ -120,10 +122,12 @@ static void fill_rand(std::vector<sycl::vec<DType, NChannels>> &v,
     }
   }();
   for (int i = 0; i < v.size(); ++i) {
-    sycl::vec<DType, NChannels> temp;
-
-    for (int j = 0; j < NChannels; j++) {
-      temp[j] = static_cast<DType>(distribution(generator));
+    T temp;
+    if constexpr (NChannels == 1) {
+      temp = static_cast<DType>(distribution(generator));
+    } else {
+      for (int j = 0; j < NChannels; ++j)
+        temp[j] = static_cast<DType>(distribution(generator));
     }
 
     v[i] = temp;

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

@@ -275,8 +275,6 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> localSize,
     imgType = VK_IMAGE_TYPE_3D;
   }
 
-  using VecType = sycl::vec<DType, NChannels>;
-
   VkFormat format = vkutil::to_vulkan_format(COrder, CType);
   const size_t imageSizeBytes = numElems * NChannels * sizeof(DType);
 
@@ -316,7 +314,7 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> localSize,
 
   printString("Populating staging buffer\n");
   // Populate staging memory
-  VecType *inputStagingData = nullptr;
+  DType *inputStagingData = nullptr;
   VK_CHECK_CALL(vkMapMemory(vk_device, inputStagingMemory, 0 /*offset*/,
                             imageSizeBytes, 0 /*flags*/,
                             (void **)&inputStagingData));
@@ -329,8 +327,9 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> localSize,
     return i;
   };
   for (int i = 0; i < numElems; ++i) {
-    inputStagingData[i] =
-        bindless_helpers::init_vector<DType, NChannels>(getInputValue(i));
+    DType v = getInputValue(i);
+    for (int j = 0; j < NChannels; ++j)
+      inputStagingData[i * NChannels + j] = v;
   }
   vkUnmapMemory(vk_device, inputStagingMemory);
 

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

@@ -303,31 +303,29 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> local_size,
 
   printString("Populating staging buffer\n");
   // Populate staging memory
-  using VecType = sycl::vec<DType, NChannels>;
-  auto init =
-      bindless_helpers::init_vector<DType, NChannels>(static_cast<DType>(0));
-
-  std::vector<VecType> input_vector_0(num_elems, init);
+  std::vector<DType> input_vector_0(num_elems * NChannels,
+                                    static_cast<DType>(0));
   std::srand(seed);
   bindless_helpers::fill_rand(input_vector_0);
 
-  VecType *inputStagingData = nullptr;
+  DType *inputStagingData = nullptr;
   VK_CHECK_CALL(vkMapMemory(vk_device, inVkImgRes1.stagingMemory, 0 /*offset*/,
                             imageSizeBytes, 0 /*flags*/,
                             (void **)&inputStagingData));
-  for (int i = 0; i < num_elems; ++i) {
+  for (int i = 0; i < (num_elems * NChannels); ++i) {
     inputStagingData[i] = input_vector_0[i];
   }
   vkUnmapMemory(vk_device, inVkImgRes1.stagingMemory);
 
-  std::vector<VecType> input_vector_1(num_elems, init);
+  std::vector<DType> input_vector_1(num_elems * NChannels,
+                                    static_cast<DType>(0));
   std::srand(seed);
   bindless_helpers::fill_rand(input_vector_1);
 
   VK_CHECK_CALL(vkMapMemory(vk_device, inVkImgRes2.stagingMemory, 0 /*offset*/,
                             imageSizeBytes, 0 /*flags*/,
                             (void **)&inputStagingData));
-  for (int i = 0; i < num_elems; ++i) {
+  for (int i = 0; i < (num_elems * NChannels); ++i) {
     inputStagingData[i] = input_vector_1[i];
   }
   vkUnmapMemory(vk_device, inVkImgRes2.stagingMemory);
@@ -535,22 +533,21 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> local_size,
   printString("Validating\n");
   // Validate that SYCL made changes to the memory
   bool validated = true;
-  VecType *outputStagingData = nullptr;
+  DType *outputStagingData = nullptr;
   VK_CHECK_CALL(vkMapMemory(vk_device, outVkImgRes.stagingMemory, 0 /*offset*/,
                             imageSizeBytes, 0 /*flags*/,
                             (void **)&outputStagingData));
-  for (int i = 0; i < num_elems; ++i) {
-    VecType expected = input_vector_0[i] + input_vector_1[i];
-    for (int j = 0; j < NChannels; ++j) {
-      // Use helper function to determine if data is accepted
-      // For integers, exact results are expected
-      // For floats, accepted error variance is passed
-      if (!util::is_equal(outputStagingData[i][j], expected[j])) {
-        std::cerr << "Result mismatch! actual[" << i << "][" << j
-                  << "] == " << outputStagingData[i][j]
-                  << " : expected == " << expected[j] << "\n";
-        validated = false;
-      }
+
+  for (int i = 0; i < (num_elems * NChannels); ++i) {
+    DType expected = input_vector_0[i] + input_vector_1[i];
+    // Use helper function to determine if data is accepted
+    // For integers, exact results are expected
+    // For floats, accepted error variance is passed
+    if (!util::is_equal(outputStagingData[i], expected)) {
+      std::cerr << "Result mismatch! actual[" << i
+                << "] == " << outputStagingData[i]
+                << " : expected == " << expected << "\n";
+      validated = false;
     }
     if (!validated)
       break;