[ESIMD] Enable math functions: cos,sin,exp,log

Author: DenisBakhvalov

SYCL/ESIMD/ext_math.cpp

@@ -16,6 +16,7 @@
 
 #include <CL/sycl.hpp>
 #include <CL/sycl/INTEL/esimd.hpp>
+#include <CL/sycl/builtins_esimd.hpp>
 #include <iostream>
 
 using namespace cl::sycl;
@@ -35,7 +36,16 @@ struct InitDataFuncWide {
 struct InitDataFuncNarrow {
   void operator()(float *In, float *Out, size_t Size) const {
     for (auto I = 0; I < Size; ++I) {
-      In[I] = 2.0f + 16.0f * ((float)I / (float)(Size - 1)); // in [2..16] range
+      In[I] = 2.0f + 16.0f * ((float)I / (float)(Size - 1)); // in [2..18] range
+      Out[I] = (float)0.0;
+    }
+  }
+};
+
+struct InitDataInRange0_5 {
+  void operator()(float *In, float *Out, size_t Size) const {
+    for (auto I = 0; I < Size; ++I) {
+      In[I] = 5.0f * ((float)I / (float)(Size - 1)); // in [0..5] range
       Out[I] = (float)0.0;
     }
   }
@@ -61,10 +71,19 @@ template <MathOp Op> float HostMathFunc(float X);
     }                                                                          \
   }
 
-DEFINE_OP(sin, sin);
-DEFINE_OP(cos, cos);
-DEFINE_OP(exp, exp);
-DEFINE_OP(log, log);
+#define DEFINE_OP_REUSE_SCALAR(Op, HostOp)                                     \
+  template <> float HostMathFunc<MathOp::Op>(float X) { return HostOp(X); }    \
+  template <int VL> struct DeviceMathFunc<VL, MathOp::Op> {                    \
+    simd<float, VL>                                                            \
+    operator()(const simd<float, VL> &X) const SYCL_ESIMD_FUNCTION {           \
+      return sycl::Op<VL>(X);                                                  \
+    }                                                                          \
+  }
+
+DEFINE_OP_REUSE_SCALAR(sin, sin);
+DEFINE_OP_REUSE_SCALAR(cos, cos);
+DEFINE_OP_REUSE_SCALAR(exp, exp);
+DEFINE_OP_REUSE_SCALAR(log, log);
 DEFINE_OP(inv, 1.0f /);
 DEFINE_OP(sqrt, sqrt);
 DEFINE_OP(rsqrt, 1.0f / sqrt);
@@ -159,13 +178,10 @@ template <int VL> bool test(queue &Q) {
   Pass &= test<MathOp::sqrt, VL>(Q, "sqrt", InitDataFuncWide{});
   Pass &= test<MathOp::inv, VL>(Q, "inv");
   Pass &= test<MathOp::rsqrt, VL>(Q, "rsqrt");
-// TODO enable these tests after the implementation is fixed
-#if ENABLE_SIN_COS_EXP_LOG
   Pass &= test<MathOp::sin, VL>(Q, "sin", InitDataFuncWide{});
   Pass &= test<MathOp::cos, VL>(Q, "cos", InitDataFuncWide{});
-  Pass &= test<MathOp::exp, VL>(Q, "exp");
+  Pass &= test<MathOp::exp, VL>(Q, "exp", InitDataInRange0_5{});
   Pass &= test<MathOp::log, VL>(Q, "log", InitDataFuncWide{});
-#endif
   return Pass;
 }