[SYCL][E2E][Bindless] Fix input value out-of-range for int8/int16 types (#15086)

globalSize.size() may be larger than maximum value of int8/int16 types.
Resolve #14945 (comment)

Author: wenju-he

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

@@ -7,6 +7,10 @@
 // Uncomment to print additional test information
 // #define VERBOSE_PRINT
 
+#ifdef _WIN32
+#define NOMINMAX
+#endif
+
 #include "../helpers/common.hpp"
 #include "vulkan_common.hpp"
 
@@ -176,11 +180,17 @@ bool run_sycl(InteropHandleT inputInteropMemHandle,
 
   printString("Validating\n");
   bool validated = true;
+  auto getExpectedValue = [&](int i) -> OutType {
+    if (CType == sycl::image_channel_type::unorm_int8)
+      return 0.5f;
+    if constexpr (std::is_integral_v<OutType>)
+      i = i % std::numeric_limits<OutType>::max();
+    return i / 2.f;
+  };
   for (int i = 0; i < globalSize.size(); i++) {
     bool mismatch = false;
     VecType expected =
-        bindless_helpers::init_vector<OutType, NChannels>(static_cast<OutType>(
-            CType == sycl::image_channel_type::unorm_int8 ? 0.5f : (i / 2.f)));
+        bindless_helpers::init_vector<OutType, NChannels>(getExpectedValue(i));
     if (!bindless_helpers::equal_vec<OutType, NChannels>(out[i], expected)) {
       mismatch = true;
       validated = false;
@@ -269,10 +279,16 @@ bool run_test(sycl::range<NDims> dims, sycl::range<NDims> localSize,
   VK_CHECK_CALL(vkMapMemory(vk_device, inputStagingMemory, 0 /*offset*/,
                             imageSizeBytes, 0 /*flags*/,
                             (void **)&inputStagingData));
+  auto getInputValue = [&](int i) -> DType {
+    if (CType == sycl::image_channel_type::unorm_int8)
+      return 255;
+    if constexpr (std::is_integral_v<DType>)
+      i = i % std::numeric_limits<DType>::max();
+    return i;
+  };
   for (int i = 0; i < numElems; ++i) {
     inputStagingData[i] =
-        bindless_helpers::init_vector<DType, NChannels>(static_cast<DType>(
-            CType == sycl::image_channel_type::unorm_int8 ? 255 : i));
+        bindless_helpers::init_vector<DType, NChannels>(getInputValue(i));
   }
   vkUnmapMemory(vk_device, inputStagingMemory);