[SYCL][ESIMD][EMU] tolerated mismatch rate in binary files comparison

SYCL/ESIMD/esimd_test_utils.hpp

@@ -16,6 +16,7 @@
 #include <chrono>
 #include <cstring>
 #include <fstream>
+#include <iomanip>
 #include <iostream>
 #include <iterator>
 #include <string>
@@ -98,28 +99,88 @@ bool write_binary_file(const char *fname, const std::vector<T> &vec,
   return !ofs.bad();
 }
 
-template <typename T>
-bool cmp_binary_files(const char *fname1, const char *fname2, T tolerance) {
-  const auto vec1 = read_binary_file<T>(fname1);
-  const auto vec2 = read_binary_file<T>(fname2);
-  if (vec1.size() != vec2.size()) {
-    std::cerr << fname1 << " size is " << vec1.size();
-    std::cerr << " whereas " << fname2 << " size is " << vec2.size()
+template <typename T,
+          typename std::enable_if<std::is_integral<T>::value ||
+                                      std::is_floating_point<T>::value,
+                                  int>::type = 0>
+bool cmp_binary_files(const char *testOutFile, const char *referenceFile,
+                      const T tolerance = 0,
+                      const double mismatchRateTolerance = 0,
+                      const int mismatchReportLimit = 9) {
+
+  if (mismatchRateTolerance) {
+    if (mismatchRateTolerance >= 1 || mismatchRateTolerance < 0) {
+      std::cerr << "Tolerated mismatch rate (" << mismatchRateTolerance
+                << ") must be set within [0, 1) range" << std::endl;
+      return false;
+    }
+
+    std::cerr << "Tolerated mismatch rate set to " << mismatchRateTolerance
               << std::endl;
+  }
+
+  const auto testVec = read_binary_file<T>(testOutFile);
+  const auto referenceVec = read_binary_file<T>(referenceFile);
+
+  if (testVec.size() != referenceVec.size()) {
+    std::cerr << testOutFile << " size is " << testVec.size();
+    std::cerr << " whereas " << referenceFile << " size is "
+              << referenceVec.size() << std::endl;
     return false;
   }
-  for (size_t i = 0; i < vec1.size(); i++) {
-    if (abs(vec1[i] - vec2[i]) > tolerance) {
-      std::cerr << "Mismatch at " << i << ' ';
-      if (sizeof(T) == 1) {
-        std::cerr << (int)vec1[i] << " vs " << (int)vec2[i] << std::endl;
-      } else {
-        std::cerr << vec1[i] << " vs " << vec2[i] << std::endl;
+
+  size_t totalMismatches = 0;
+  const size_t size = testVec.size();
+  double maxRelativeDiff = 0;
+  bool status = true;
+  for (size_t i = 0; i < size; i++) {
+    const auto diff = abs(testVec[i] - referenceVec[i]);
+    if (diff > tolerance) {
+      if (!mismatchRateTolerance || (totalMismatches < mismatchReportLimit)) {
+
+        std::cerr << "Mismatch at " << i << ' ';
+        if (sizeof(T) == 1) {
+          std::cerr << (int)testVec[i] << " vs " << (int)referenceVec[i];
+        } else {
+          std::cerr << testVec[i] << " vs " << referenceVec[i];
+        }
+
+        maxRelativeDiff = std::max(maxRelativeDiff,
+                                   static_cast<double>(diff) / referenceVec[i]);
+
+        if (!mismatchRateTolerance) {
+          std::cerr << std::endl;
+          status = false;
+          break;
+        } else {
+          std::cerr << ". Current mismatch rate: " << std::setprecision(8)
+                    << std::fixed << static_cast<double>(totalMismatches) / size
+                    << std::endl;
+        }
+
+      } else if (totalMismatches == mismatchReportLimit) {
+        std::cerr << "Mismatch output stopped." << std::endl;
       }
-      return false;
+
+      totalMismatches++;
+    }
+  }
+
+  if (totalMismatches) {
+    const auto totalMismatchRate = static_cast<double>(totalMismatches) / size;
+    if (totalMismatchRate > mismatchRateTolerance) {
+      std::cerr << "Mismatch rate of " << totalMismatchRate
+                << " has exceeded the tolerated amount of "
+                << mismatchRateTolerance << std::endl;
+      status = false;
     }
+
+    std::cerr << "Total mismatch rate is " << totalMismatchRate
+              << " with max relative difference of " << maxRelativeDiff
+              << std::endl;
   }
-  return true;
+
+  return status;
 }
 
 // dump every element of sequence [first, last) to std::cout

SYCL/ESIMD/mandelbrot/mandelbrot.cpp

@@ -36,6 +36,10 @@ using namespace sycl::ext::intel::esimd;
 #define WIDTH 800
 #define HEIGHT 602
 
+#define EMU_TOTAL_MISMATCH_RATE_TOLERANCE                                      \
+  0.0027 // the observed total mismatch rate (due to inherent divergency in host
+         // vs GPU FP computations)
+
 template <typename ACC>
 ESIMD_INLINE void mandelbrot(ACC out_image, int ix, int iy, int crunch,
                              float xOff, float yOff, float scale) {
@@ -94,6 +98,9 @@ int main(int argc, char *argv[]) {
   double kernel_times = 0;
   unsigned num_iters = 10;
 
+  queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler(),
+          property::queue::enable_profiling{});
+
   try {
     cl::sycl::image<2> imgOutput((unsigned int *)buf, image_channel_order::rgba,
                                  image_channel_type::unsigned_int8,
@@ -107,9 +114,6 @@ int main(int argc, char *argv[]) {
     // Number of workitems in a workgroup
     cl::sycl::range<2> LocalRange{1, 1};
 
-    queue q(esimd_test::ESIMDSelector{}, esimd_test::createExceptionHandler(),
-            property::queue::enable_profiling{});
-
     auto dev = q.get_device();
     auto ctxt = q.get_context();
     std::cout << "Running on " << dev.get_info<info::device::name>() << "\n";
@@ -164,7 +168,11 @@ int main(int argc, char *argv[]) {
   fclose(dumpfile);
 
   bool passed = true;
-  if (!esimd_test::cmp_binary_files<unsigned char>(out_file, argv[2], 0)) {
+  if (!esimd_test::cmp_binary_files<unsigned char>(
+          out_file, argv[2], 0,
+          q.get_backend() == cl::sycl::backend::ext_intel_esimd_emulator
+              ? EMU_TOTAL_MISMATCH_RATE_TOLERANCE
+              : 0)) {
     std::cerr << out_file << " does not match the reference file " << argv[2]
               << std::endl;
     passed = false;