Merge branch 'unisa-hpc:master' into sycl-mlir

pattern/reduction.cpp

@@ -61,8 +61,7 @@ class Reduction
   }
 
   bool verify(VerificationSetting &ver) {
-    T result = _final_output_buff->template get_access<sycl::access::mode::read>(
-        sycl::range<1>{0}, sycl::id<1>{1})[0];
+    T result = _final_output_buff->get_host_access()[0];
 
     // Calculate CPU result in fp64 to avoid obtaining a wrong verification result
     std::vector<double> input_fp64(_input.size());
@@ -71,7 +70,7 @@ class Reduction
 
     double delta =
         static_cast<double>(result) - std::accumulate(input_fp64.begin(), input_fp64.end(), T{});
-    
+
     return std::abs(delta) < 1.e-5;
   }
 private:

single-kernel/nbody.cpp

@@ -2,6 +2,7 @@
 
 #include <iostream>
 #include <cassert>
+#include <limits>
 
 using namespace cl;
 
@@ -105,41 +106,19 @@ class NBody
       host_resulting_velocities[i] = new_v;
     }
 
-    double deviation = std::sqrt(
-        calculateSquaredDifference(host_resulting_particles.data(), resulting_particles.get_pointer(), particles.size()) +
-        calculateSquaredDifference(host_resulting_velocities.data(), resulting_velocities.get_pointer(), particles.size()));
-
-    return deviation < 1.e-6;
+    constexpr float_type maxErr = 10.f * std::numeric_limits<float_type>::epsilon();
+    return checkResults(host_resulting_particles.begin(), host_resulting_particles.end(),
+               resulting_particles.get_pointer(), maxErr) &&
+           checkResults(host_resulting_velocities.begin(), host_resulting_velocities.end(),
+               resulting_velocities.get_pointer(), maxErr);
   }
 
 protected:
-
-  template<class T>
-  double calculateSquaredDifference(sycl::vec<T,3> a, sycl::vec<T,3> b) {
-    auto diff = a - b;
-    diff *= diff;
-
-    return static_cast<float_type>(diff.x()+diff.y()+diff.z());
-  }
-
-  template<class T>
-  double calculateSquaredDifference(sycl::vec<T,4> a, sycl::vec<T,4> b) {
-    auto diff = a - b;
-    diff *= diff;
-
-    return static_cast<float_type>(diff.x()+diff.y()+diff.z()+diff.w());
-  }
-
-  template<class T>
-  double calculateSquaredDifference(const T* a, const T* b, std::size_t size) {
-
-    double result = 0.0;
-
-    for(std::size_t i = 0; i < size; ++i) {
-      result += calculateSquaredDifference(a[i], b[i]);
-    }
-
-    return result;
+  template <class InputIter0, class InputIter1>
+  static bool checkResults(InputIter0 expectedBegin, InputIter0 expectedEnd, InputIter1 gotBegin, float_type maxErr) {
+    return std::equal(expectedBegin, expectedEnd, gotBegin, [=](const auto& expected, const auto& got) {
+      return sycl::distance(expected, got) / sycl::length(expected) < maxErr;
+    });
   }
 
   void submitNDRange(sycl::buffer<particle_type>& particles, sycl::buffer<vector_type>& velocities) {