Merge pull request #892 from steffenlarsen/steffen/add_absolute_toler…

…ance_accuracy_mode

Add "absolute error tolerance" accuracy mode check

tests/math_builtin_api/math_builtin.h

@@ -79,16 +79,27 @@ inline sycl::half min_t<sycl::half>() {
   return static_cast<sycl::half>(powf(2.0f, -14.0f));
 }
 
+enum class AccuracyMode { ULP, AbsoluteTolerance };
+
+inline std::string GetAccuracyModeStr(AccuracyMode accuracy_mode) {
+  switch (accuracy_mode) {
+    case AccuracyMode::ULP:
+      return "ULP";
+    case AccuracyMode::AbsoluteTolerance:
+      return "absolute error tolerance";
+  }
+}
+
 template <typename T>
-bool verify(sycl_cts::util::logger& log, T a, T b, int accuracy,
-            const std::string& comment);
+bool verify(sycl_cts::util::logger& log, T a, T b, float accuracy,
+            AccuracyMode accuracy_mode, const std::string& comment);
 
 template <typename T>
 typename std::enable_if<std::is_floating_point<T>::value ||
                             std::is_same<sycl::half, T>::value,
                         bool>::type
 verify(sycl_cts::util::logger& log, T value, sycl_cts::resultRef<T> r,
-       int accuracy, const std::string& comment) {
+       float accuracy, AccuracyMode accuracy_mode, const std::string& comment) {
   const T reference = r.res;
 
   if (!r.undefined.empty())
@@ -101,30 +112,42 @@ verify(sycl_cts::util::logger& log, T value, sycl_cts::resultRef<T> r,
   if (!std::isnan(value) && !std::isnan(reference) && !std::isinf(value) &&
       !std::isinf(reference)) {
     if (accuracy < 0)
-      return true;  // Implementation-defined or infinite ULP according to spec
+      return true;  // Implementation-defined or infinite ULP/tolerance
+                    // according to spec
 
     if ((std::fabs(value) < min_t<T>()) && (std::fabs(reference) < min_t<T>()))
       return true;  // Subnormal numbers are the lower border for comparison
 
-    const auto ulpsExpected = static_cast<unsigned int>(accuracy);
     const T difference = static_cast<T>(std::fabs(value - reference));
-    const T differenceExpected = ulpsExpected * get_ulp_std(reference);
-
-    if (difference <= differenceExpected) return true;
+    switch (accuracy_mode) {
+      case AccuracyMode::ULP: {
+        const auto ulpsExpected = static_cast<unsigned int>(accuracy);
+        const T differenceExpected = ulpsExpected * get_ulp_std(reference);
+
+        if (difference <= differenceExpected) return true;
+        break;
+      }
+      case AccuracyMode::AbsoluteTolerance: {
+        if (difference <= accuracy) return true;
+        break;
+      }
+    }
   }
 
   log.note("value: " + printable(value) +
            ", reference: " + printable(reference));
-  std::string msg = "Expected accuracy in ULP: " + std::to_string(accuracy);
+  std::string msg = "Expected accuracy in " +
+                    GetAccuracyModeStr(accuracy_mode) + ": " +
+                    std::to_string(accuracy);
   if (!comment.empty()) msg += ", " + comment;
   log.note(msg);
   return false;
 }
 
 template <typename T>
 typename std::enable_if_t<std::is_integral_v<T>, bool> verify(
-    sycl_cts::util::logger& log, T value, sycl_cts::resultRef<T> r, int,
-    const std::string&) {
+    sycl_cts::util::logger& log, T value, sycl_cts::resultRef<T> r, float,
+    AccuracyMode, const std::string&) {
   bool result = value == r.res || !r.undefined.empty();
   if (!result)
     log.note("value: " + std::to_string(value) +
@@ -134,37 +157,39 @@ typename std::enable_if_t<std::is_integral_v<T>, bool> verify(
 
 template <typename T, int N>
 bool verify(sycl_cts::util::logger& log, sycl::vec<T, N> a,
-            sycl_cts::resultRef<sycl::vec<T, N>> r, int accuracy,
-            const std::string& comment) {
+            sycl_cts::resultRef<sycl::vec<T, N>> r, float accuracy,
+            AccuracyMode accuracy_mode, const std::string& comment) {
   sycl::vec<T, N> b = r.res;
   for (int i = 0; i < sycl_cts::math::numElements(a); i++)
     if (r.undefined.find(i) == r.undefined.end() &&
-        !verify(log, a[i], b[i], accuracy, comment))
+        !verify(log, a[i], b[i], accuracy, accuracy_mode, comment))
       return false;
   return true;
 }
 
 template <typename T, size_t N>
 bool verify(sycl_cts::util::logger& log, sycl::marray<T, N> a,
-            sycl_cts::resultRef<sycl::marray<T, N>> r, int accuracy,
-            const std::string& comment) {
+            sycl_cts::resultRef<sycl::marray<T, N>> r, float accuracy,
+            AccuracyMode accuracy_mode, const std::string& comment) {
   sycl::marray<T, N> b = r.res;
   for (size_t i = 0; i < N; i++)
     if (r.undefined.find(i) == r.undefined.end() &&
-        !verify(log, a[i], b[i], accuracy, comment))
+        !verify(log, a[i], b[i], accuracy, accuracy_mode, comment))
       return false;
   return true;
 }
 
 template <typename T>
-bool verify(sycl_cts::util::logger& log, T a, T b, int accuracy,
-            const std::string& comment) {
-  return verify(log, a, sycl_cts::resultRef<T>(b), accuracy, comment);
+bool verify(sycl_cts::util::logger& log, T a, T b, float accuracy,
+            AccuracyMode accuracy_mode, const std::string& comment) {
+  return verify(log, a, sycl_cts::resultRef<T>(b), accuracy, accuracy_mode,
+                comment);
 }
 
 template <int N, typename returnT, typename funT>
 void check_function(sycl_cts::util::logger& log, funT fun,
-                    sycl_cts::resultRef<returnT> ref, int accuracy = 0,
+                    sycl_cts::resultRef<returnT> ref, float accuracy = 0.0f,
+                    AccuracyMode accuracy_mode = AccuracyMode::ULP,
                     const std::string& comment = {}) {
   sycl::range<1> ndRng(1);
   returnT kernelResult;
@@ -183,7 +208,7 @@ void check_function(sycl_cts::util::logger& log, funT fun,
     FAIL(log, errorMsg.c_str());
   }
 
-  if (!verify(log, kernelResult, ref, accuracy, comment))
+  if (!verify(log, kernelResult, ref, accuracy, accuracy_mode, comment))
     FAIL(log,
          "tests case: " + std::to_string(N) + ". Correctness check failed.");
 
@@ -193,14 +218,15 @@ void check_function(sycl_cts::util::logger& log, funT fun,
        ". Correctness check failed on host.");
   // SYCL 2020 specification sets no requirements for math built-ins accuracy
   // on host, hence passing negative value to 'verify' helper to indicate that.
-  CHECK(verify(log, hostRes, ref, -1, comment));
+  CHECK(verify(log, hostRes, ref, -1, accuracy_mode, comment));
 }
 
 template <int N, typename returnT, typename funT, typename argT>
-void check_function_multi_ptr_private(sycl_cts::util::logger& log, funT fun,
-                                      sycl_cts::resultRef<returnT> ref,
-                                      argT ptrRef, int accuracy = 0,
-                                      const std::string& comment = {}) {
+void check_function_multi_ptr_private(
+    sycl_cts::util::logger& log, funT fun, sycl_cts::resultRef<returnT> ref,
+    argT ptrRef, float accuracy = 0.0f,
+    AccuracyMode accuracy_mode = AccuracyMode::ULP,
+    const std::string& comment = {}) {
   sycl::range<1> ndRng(1);
   returnT kernelResult;
   argT kernelResultArg;
@@ -225,10 +251,10 @@ void check_function_multi_ptr_private(sycl_cts::util::logger& log, funT fun,
     FAIL(log, errorMsg.c_str());
   }
 
-  if (!verify(log, kernelResult, ref, accuracy, comment))
+  if (!verify(log, kernelResult, ref, accuracy, accuracy_mode, comment))
     FAIL(log,
          "tests case: " + std::to_string(N) + ". Correctness check failed.");
-  if (!verify(log, kernelResultArg, ptrRef, accuracy, comment))
+  if (!verify(log, kernelResultArg, ptrRef, accuracy, accuracy_mode, comment))
     FAIL(log, "tests case: " + std::to_string(N) +
                   ". Correctness check for ptr failed.");
 
@@ -237,20 +263,22 @@ void check_function_multi_ptr_private(sycl_cts::util::logger& log, funT fun,
   {
     INFO("tests case: " + std::to_string(N) +
          ". Correctness check failed on host.");
-    CHECK(verify(log, hostRes.res, ref, accuracy, comment));
+    CHECK(verify(log, hostRes.res, ref, accuracy, accuracy_mode, comment));
   }
   {
     INFO("tests case: " + std::to_string(N) +
          ". Correctness check for ptr failed on host.");
-    CHECK(verify(log, hostRes.resArg, ptrRef, accuracy, comment));
+    CHECK(
+        verify(log, hostRes.resArg, ptrRef, accuracy, accuracy_mode, comment));
   }
 }
 
 template <int N, typename returnT, typename funT, typename argT>
-void check_function_multi_ptr_global(sycl_cts::util::logger& log, funT fun,
-                                     argT arg, sycl_cts::resultRef<returnT> ref,
-                                     argT ptrRef, int accuracy = 0,
-                                     const std::string& comment = {}) {
+void check_function_multi_ptr_global(
+    sycl_cts::util::logger& log, funT fun, argT arg,
+    sycl_cts::resultRef<returnT> ref, argT ptrRef, float accuracy = 0.0f,
+    AccuracyMode accuracy_mode = AccuracyMode::ULP,
+    const std::string& comment = {}) {
   sycl::range<1> ndRng(1);
   returnT kernelResult;
   auto&& testQueue = once_per_unit::get_queue();
@@ -271,19 +299,20 @@ void check_function_multi_ptr_global(sycl_cts::util::logger& log, funT fun,
     FAIL(log, errorMsg.c_str());
   }
 
-  if (!verify(log, kernelResult, ref, accuracy, comment))
+  if (!verify(log, kernelResult, ref, accuracy, accuracy_mode, comment))
     FAIL(log,
          "tests case: " + std::to_string(N) + ". Correctness check failed.");
-  if (!verify(log, arg, ptrRef, accuracy, comment))
+  if (!verify(log, arg, ptrRef, accuracy, accuracy_mode, comment))
     FAIL(log, "tests case: " + std::to_string(N) +
                   ". Correctness check for ptr failed.");
 }
 
 template <int N, typename returnT, typename funT, typename argT>
-void check_function_multi_ptr_local(sycl_cts::util::logger& log, funT fun,
-                                    argT arg, sycl_cts::resultRef<returnT> ref,
-                                    argT ptrRef, int accuracy = 0,
-                                    const std::string& comment = {}) {
+void check_function_multi_ptr_local(
+    sycl_cts::util::logger& log, funT fun, argT arg,
+    sycl_cts::resultRef<returnT> ref, argT ptrRef, float accuracy = 0.0f,
+    AccuracyMode accuracy_mode = AccuracyMode::ULP,
+    const std::string& comment = {}) {
   sycl::range<1> ndRng(1);
   returnT kernelResult;
   auto&& testQueue = once_per_unit::get_queue();
@@ -310,10 +339,10 @@ void check_function_multi_ptr_local(sycl_cts::util::logger& log, funT fun,
     FAIL(log, errorMsg.c_str());
   }
 
-  if (!verify(log, kernelResult, ref, accuracy, comment))
+  if (!verify(log, kernelResult, ref, accuracy, accuracy_mode, comment))
     FAIL(log,
          "tests case: " + std::to_string(N) + ". Correctness check failed.");
-  if (!verify(log, arg, ptrRef, accuracy, comment))
+  if (!verify(log, arg, ptrRef, accuracy, accuracy_mode, comment))
     FAIL(log, "tests case: " + std::to_string(N) +
                   ". Correctness check for ptr failed.");
 }

tests/math_builtin_api/modules/sycl_functions.py

@@ -1,6 +1,6 @@
 """Represents a function signature."""
 class funsig:
-    def __init__(self, namespace, ret_type, name, arg_types=[], accuracy="", comment="", pntr_indx=[], mutations=[], template_arg_map=[]):
+    def __init__(self, namespace, ret_type, name, arg_types=[], accuracy="", comment="", pntr_indx=[], mutations=[], template_arg_map=[], accuracy_mode="ULP"):
         self.namespace = namespace # Namespace of function.
         self.ret_type = ret_type # Function return type.
         self.name = name # Function name.
@@ -18,6 +18,7 @@ def __init__(self, namespace, ret_type, name, arg_types=[], accuracy="", comment
         # e.g. int32_t and float are OK, but not int64_t and marray<float, 2>)
         self.template_arg_map = template_arg_map # List of indices mapping template arugments to function argument types.
         # An empty list signifies a non-templated function.
+        self.accuracy_mode = accuracy_mode # A mode for accuracy expectations for a given builtin. Must be either "ULP" or "AbsoluteTolerance"
     def __eq__(self, other):
         if isinstance(other, funsig):
             return ((self.namespace == other.namespace) and
@@ -28,13 +29,14 @@ def __eq__(self, other):
                     (self.comment == other.comment) and
                     (self.pntr_indx == other.pntr_indx) and
                     (self.mutations == other.mutations) and
-                    (self.template_arg_map == other.template_arg_map))
+                    (self.template_arg_map == other.template_arg_map) and
+                    (self.accuracy_mode == other.accuracy_mode))
         else:
             return False
     def __ne__(self, other):
         return (not self.__eq__(other))
     def __hash__(self):
-        return hash((self.namespace, self.ret_type, self.name, str(self.arg_types), self.accuracy, self.comment, str(self.pntr_indx), str(self.mutations), str(self.template_arg_map)))
+        return hash((self.namespace, self.ret_type, self.name, str(self.arg_types), self.accuracy, self.comment, str(self.pntr_indx), str(self.mutations), str(self.template_arg_map), self.accuracy_mode))
 
 def create_integer_signatures():
     sig_list = []
@@ -164,16 +166,16 @@ def create_common_signatures():
     f_min_4 = funsig("sycl", "genfloath", "min", ["genfloath", "sycl::half"], template_arg_map=[0])
     sig_list.append(f_min_4)
 
-    f_mix = funsig("sycl", "genfloat", "mix", ["genfloat", "genfloat", "genfloat"], "1", template_arg_map=[0,1,2])
+    f_mix = funsig("sycl", "genfloat", "mix", ["genfloat", "genfloat", "genfloat"], "1e-3", template_arg_map=[0,1,2], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_mix)
 
-    f_mix_2 = funsig("sycl", "genfloatf", "mix", ["genfloatf", "genfloatf", "float"], "1", template_arg_map=[0,1])
+    f_mix_2 = funsig("sycl", "genfloatf", "mix", ["genfloatf", "genfloatf", "float"], "1e-3", template_arg_map=[0,1], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_mix_2)
 
-    f_mix_3 = funsig("sycl", "genfloatd", "mix", ["genfloatd", "genfloatd", "double"], "1", template_arg_map=[0,1])
+    f_mix_3 = funsig("sycl", "genfloatd", "mix", ["genfloatd", "genfloatd", "double"], "1e-3", template_arg_map=[0,1], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_mix_3)
 
-    f_mix_4 = funsig("sycl", "genfloath", "mix", ["genfloath", "genfloath", "sycl::half"], "1", template_arg_map=[0,1])
+    f_mix_4 = funsig("sycl", "genfloath", "mix", ["genfloath", "genfloath", "sycl::half"], "1e-3", template_arg_map=[0,1], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_mix_4)
 
     f_radians = funsig("sycl", "genfloat", "radians", ["genfloat"], "3", template_arg_map=[0])
@@ -191,16 +193,16 @@ def create_common_signatures():
     f_step_4 = funsig("sycl", "genfloath", "step", ["sycl::half", "genfloath"], template_arg_map=[1])
     sig_list.append(f_step_4)
 
-    f_smoothstep = funsig("sycl", "genfloat", "smoothstep", ["genfloat", "genfloat", "genfloat"], template_arg_map=[0,1,2])
+    f_smoothstep = funsig("sycl", "genfloat", "smoothstep", ["genfloat", "genfloat", "genfloat"], "1e-5", template_arg_map=[0,1,2], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_smoothstep)
 
-    f_smoothstep_2 = funsig("sycl", "genfloatf", "smoothstep", ["float", "float", "genfloatf"], template_arg_map=[2])
+    f_smoothstep_2 = funsig("sycl", "genfloatf", "smoothstep", ["float", "float", "genfloatf"], "1e-5", template_arg_map=[2], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_smoothstep_2)
 
-    f_smoothstep_3 = funsig("sycl", "genfloatd", "smoothstep", ["double", "double", "genfloatd"], template_arg_map=[2])
+    f_smoothstep_3 = funsig("sycl", "genfloatd", "smoothstep", ["double", "double", "genfloatd"], "1e-5", template_arg_map=[2], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_smoothstep_3)
 
-    f_smoothstep_4 = funsig("sycl", "genfloath", "smoothstep", ["sycl::half", "sycl::half", "genfloath"], template_arg_map=[2])
+    f_smoothstep_4 = funsig("sycl", "genfloath", "smoothstep", ["sycl::half", "sycl::half", "genfloath"], "1e-5", template_arg_map=[2], accuracy_mode="AbsoluteTolerance")
     sig_list.append(f_smoothstep_4)
 
     f_sign = funsig("sycl", "genfloat", "sign", ["genfloat"], template_arg_map=[0])

tests/math_builtin_api/modules/test_generator.py

@@ -317,11 +317,12 @@ def generate_test_case(test_id, types, sig, memory, check, decorated = ""):
     testCaseSource = testCaseSource.replace("$RETURN_TYPE", sig.ret_type.name)
     if sig.accuracy:##If the signature contains an accuracy value
         accuracy = sig.accuracy
+        accuracy_mode = sig.accuracy_mode # Accuracy mode should always be set.
         # if accuracy depends on vecSize
         if "vecSize" in accuracy:
             vecSize = str(sig.arg_types[0].dim)
             accuracy = accuracy.replace("vecSize", vecSize)
-        testCaseSource = testCaseSource.replace("$ACCURACY", ", " + accuracy)
+        testCaseSource = testCaseSource.replace("$ACCURACY", ", " + accuracy + ", AccuracyMode::" + accuracy_mode)
     else:
         testCaseSource = testCaseSource.replace("$ACCURACY", "")
     if sig.comment:##If the signature contains comment for accuracy
@@ -478,7 +479,8 @@ def expand_signature(types, signature):
                                         signature.name, [matched_typelists[signature.arg_types[j]][i] 
                                                          for j in range(len(signature.arg_types))],
                                         signature.accuracy, signature.comment, signature.pntr_indx[:],
-                                        signature.mutations[:], signature.template_arg_map[:])
+                                        signature.mutations[:], signature.template_arg_map[:],
+                                        signature.accuracy_mode)
         exp_sig.append(new_sig)
 
     return exp_sig