[SYCL][E2E] Add sycl::half_precision math functions to marray_math.cpp (intel#12099)

Add missing math functions defined in the sycl::half_precision
namespace. Thank you for your review.

---------

Co-authored-by: Jin Z <5zj@equinox.ftpn.ornl.gov>
Co-authored-by: Jin Z <5zj@cousteau.ftpn.ornl.gov>

Author: 3 people

sycl/test-e2e/Basic/built-ins/marray_math.cpp

@@ -5,8 +5,37 @@
 // RUN: %if preview-breaking-changes-supported %{ %{build} %{mathflags} -fpreview-breaking-changes -o %t_preview.out %}
 // RUN: %if preview-breaking-changes-supported %{ %{run} %t_preview.out%}
 
+#include <cmath>
 #include <sycl/sycl.hpp>
 
+// Reference
+// https://github.com/KhronosGroup/SYCL-CTS/blob/SYCL-2020/util/accuracy.h
+template <typename T> T get_ulp_sycl(T x) {
+  const T inf = std::numeric_limits<T>::infinity();
+  const T negative = sycl::fabs(sycl::nextafter(x, -inf) - x);
+  const T positive = sycl::fabs(sycl::nextafter(x, inf) - x);
+  return sycl::fmin(negative, positive);
+}
+
+// Case 1: According to the SYCL 2020 specification, the sycl::half_precision
+// math functions all take float or either an marray, vec, or swizzled vec of
+// float. Additionally, the ULP expectation of all of these should be 8192
+// according to the spec.
+//
+// Case 2: The allowed error in ULP is less than 8192 for half precision
+// floating-point (sycl::half)
+template <typename T> T get_8192ulp_sycl(T x) {
+  const auto ulp = get_ulp_sycl<float>(x);
+  const float multiplier = 8192.0f;
+  return static_cast<T>(ulp * multiplier);
+}
+
+template <typename T> bool compare_floats_8192ulp(T actual, T expected) {
+  const T difference = static_cast<T>(std::fabs(actual - expected));
+  const T difference_expected = get_8192ulp_sycl(expected);
+  return difference <= difference_expected;
+}
+
 #define TEST(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, DELTA, ...)                \
   {                                                                            \
     {                                                                          \
@@ -76,6 +105,26 @@
     }                                                                          \
   }
 
+#define TEST4(FUNC, MARRAY_ELEM_TYPE, DIM, EXPECTED, ...)                      \
+  {                                                                            \
+    {                                                                          \
+      MARRAY_ELEM_TYPE result[DIM];                                            \
+      {                                                                        \
+        sycl::buffer<MARRAY_ELEM_TYPE> b(result, sycl::range{DIM});            \
+        deviceQueue.submit([&](sycl::handler &cgh) {                           \
+          sycl::accessor res_access{b, cgh};                                   \
+          cgh.single_task([=]() {                                              \
+            sycl::marray<MARRAY_ELEM_TYPE, DIM> res = FUNC(__VA_ARGS__);       \
+            for (int i = 0; i < DIM; i++)                                      \
+              res_access[i] = res[i];                                          \
+          });                                                                  \
+        });                                                                    \
+      }                                                                        \
+      for (int i = 0; i < DIM; i++)                                            \
+        assert(compare_floats_8192ulp(result[i], EXPECTED[i]));                \
+    }                                                                          \
+  }
+
 #define EXPECTED(TYPE, ...) ((TYPE[]){__VA_ARGS__})
 
 int main() {
@@ -117,8 +166,32 @@ int main() {
   TEST3(sycl::nan, float, 3, ma7);
   if (deviceQueue.get_device().has(sycl::aspect::fp64))
     TEST3(sycl::nan, double, 3, ma8);
-  TEST(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), 0.1,
-       ma1);
+
+  TEST4(sycl::half_precision::sin, float, 3,
+        EXPECTED(float, 0.98545f, -0.871576f, -0.832267f), ma6);
+  TEST4(sycl::half_precision::cos, float, 3,
+        EXPECTED(float, 0.169967f, -0.490261f, 0.554375f), ma6);
+  TEST4(sycl::half_precision::tan, float, 3,
+        EXPECTED(float, 5.797f, 1.777f, -1.501f), ma6);
+  TEST4(sycl::half_precision::divide, float, 2, EXPECTED(float, 3.0f, 1.0f),
+        ma2, ma1);
+  TEST4(sycl::half_precision::log, float, 2, EXPECTED(float, 0.0f, 0.693f),
+        ma1);
+  TEST4(sycl::half_precision::log2, float, 2, EXPECTED(float, 0.0f, 1.f), ma1);
+  TEST4(sycl::half_precision::log10, float, 2, EXPECTED(float, 0.0f, 0.301f),
+        ma1);
+  TEST4(sycl::half_precision::powr, float, 2, EXPECTED(float, 3.0f, 4.0f), ma2,
+        ma1);
+  TEST4(sycl::half_precision::recip, float, 2, EXPECTED(float, 0.333333f, 0.5f),
+        ma2);
+  TEST4(sycl::half_precision::sqrt, float, 2, EXPECTED(float, 1.0f, 1.414f),
+        ma1);
+  TEST4(sycl::half_precision::rsqrt, float, 2, EXPECTED(float, 1.0f, 0.707f),
+        ma1);
+  TEST4(sycl::half_precision::exp, float, 2, EXPECTED(float, 2.718f, 7.389f),
+        ma1);
+  TEST4(sycl::half_precision::exp2, float, 2, EXPECTED(float, 2, 4), ma1);
+  TEST4(sycl::half_precision::exp10, float, 2, EXPECTED(float, 10, 100), ma1);
 
   return 0;
 }