Add a dedicated kernel for sycl::half input dtype

Add a comment to casting in ErfFunctor implementation

Author: antonwolfy

dpnp/backend/extensions/ufunc/elementwise_functions/erf.cpp

@@ -62,10 +62,15 @@ namespace td_ns = dpctl::tensor::type_dispatch;
 template <typename T>
 struct OutputType
 {
-    using value_type =
-        typename std::disjunction<td_ns::TypeMapResultEntry<T, float>,
-                                  td_ns::TypeMapResultEntry<T, double>,
-                                  td_ns::DefaultResultEntry<void>>::result_type;
+    /**
+     * scipy>=1.16 assumes a pair 'e->d', but dpnp 'e->f' without an extra
+     * kernel 'e->d' (when fp64 supported) to reduce memory footprint
+     */
+    using value_type = typename std::disjunction<
+        td_ns::TypeMapResultEntry<T, sycl::half, float>,
+        td_ns::TypeMapResultEntry<T, float>,
+        td_ns::TypeMapResultEntry<T, double>,
+        td_ns::DefaultResultEntry<void>>::result_type;
 };
 
 using dpnp::kernels::erf::ErfFunctor;

dpnp/backend/kernels/elementwise_functions/erf.hpp

@@ -45,7 +45,14 @@ struct ErfFunctor
 
     Tp operator()(const argT &x) const
     {
-        return sycl::erf(x);
+        if constexpr (std::is_same_v<argT, sycl::half> &&
+                      std::is_same_v<Tp, float>) {
+            // cast sycl::half to float for accuracy reasons
+            return sycl::erf(float(x));
+        }
+        else {
+            return sycl::erf(x);
+        }
     }
 };
 } // namespace dpnp::kernels::erf