[SYCL][ESIMD][EMU] tolerated mismatch rate in binary files comparison (e.g. for cases resulted from deviations in GPU vs host-based FP computations).

Author: lsatanov

SYCL/ESIMD/esimd_test_utils.hpp

@@ -16,6 +16,7 @@
 #include <chrono>
 #include <cstring>
 #include <fstream>
+#include <iomanip>
 #include <iostream>
 #include <iterator>
 #include <string>
@@ -99,26 +100,73 @@ bool write_binary_file(const char *fname, const std::vector<T> &vec,
 }
 
 template <typename T>
-bool cmp_binary_files(const char *fname1, const char *fname2, T tolerance) {
+bool cmp_binary_files(const char *fname1, const char *fname2, const T tolerance,
+                      const double toleratedMismatchRate = 0,
+                      const int mismatchReportThrottleLimit = 5) {
+
+  if (toleratedMismatchRate) {
+    if (toleratedMismatchRate >= 1 || toleratedMismatchRate < 0) {
+      std::cerr << "Tolerated mismatch rate (" << toleratedMismatchRate
+                << ") must be set within [0, 1) range" << std::endl;
+      return false;
+    }
+
+    std::cerr << "Tolerated mismatch rate set to " << toleratedMismatchRate
+              << std::endl;
+  }
+
   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()
               << std::endl;
     return false;
   }
-  for (size_t i = 0; i < vec1.size(); i++) {
+
+  double totalMismatches = 0;
+  const double size = vec1.size();
+  for (size_t i = 0; i < 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;
+      if (!toleratedMismatchRate ||
+          (totalMismatches < mismatchReportThrottleLimit)) {
+
+        std::cerr << "Mismatch at " << i << ' ';
+        if (sizeof(T) == 1) {
+          std::cerr << (int)vec1[i] << " vs " << (int)vec2[i];
+        } else {
+          std::cerr << vec1[i] << " vs " << vec2[i];
+        }
+
+        if (!toleratedMismatchRate) {
+          std::cerr << std::endl;
+          return false;
+        } else {
+          std::cerr << ". Current mismatch rate: " << std::setprecision(8)
+                    << std::fixed << totalMismatches / size << std::endl;
+        }
+
+      } else if (totalMismatches == mismatchReportThrottleLimit) {
+        std::cerr << "Mismatch output throttled ... " << std::endl;
       }
+
+      totalMismatches++;
+    }
+  }
+
+  if (totalMismatches) {
+    const auto totalMismatchRate = totalMismatches / size;
+    if (totalMismatchRate > toleratedMismatchRate) {
+      std::cerr << "Mismatch rate of " << totalMismatchRate
+                << " has exceeded the tolerated amount of "
+                << toleratedMismatchRate << std::endl;
       return false;
     }
+
+    std::cerr << "Total mismatch rate is " << totalMismatchRate << std::endl;
   }
+
   return true;
 }
 

SYCL/ESIMD/mandelbrot/mandelbrot.cpp

@@ -36,6 +36,10 @@ using namespace sycl::ext::intel::experimental::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;