Use binary function template for non_equal

Author: vlad-perevezentsev

dpctl/tensor/libtensor/include/kernels/elementwise_functions/not_equal.hpp

@@ -167,32 +167,10 @@ sycl::event not_equal_contig_impl(sycl::queue exec_q,
                                   py::ssize_t res_offset,
                                   const std::vector<sycl::event> &depends = {})
 {
-    sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
-        cgh.depends_on(depends);
-
-        size_t lws = 64;
-        constexpr unsigned int vec_sz = 4;
-        constexpr unsigned int n_vecs = 2;
-        const size_t n_groups =
-            ((nelems + lws * n_vecs * vec_sz - 1) / (lws * n_vecs * vec_sz));
-        const auto gws_range = sycl::range<1>(n_groups * lws);
-        const auto lws_range = sycl::range<1>(lws);
-
-        using resTy = typename NotEqualOutputType<argTy1, argTy2>::value_type;
-
-        const argTy1 *arg1_tp =
-            reinterpret_cast<const argTy1 *>(arg1_p) + arg1_offset;
-        const argTy2 *arg2_tp =
-            reinterpret_cast<const argTy2 *>(arg2_p) + arg2_offset;
-        resTy *res_tp = reinterpret_cast<resTy *>(res_p) + res_offset;
-
-        cgh.parallel_for<
-            not_equal_contig_kernel<argTy1, argTy2, resTy, vec_sz, n_vecs>>(
-            sycl::nd_range<1>(gws_range, lws_range),
-            NotEqualContigFunctor<argTy1, argTy2, resTy, vec_sz, n_vecs>(
-                arg1_tp, arg2_tp, res_tp, nelems));
-    });
-    return comp_ev;
+    return elementwise_common::binary_contig_impl<
+        argTy1, argTy2, NotEqualOutputType, NotEqualContigFunctor,
+        not_equal_contig_kernel>(exec_q, nelems, arg1_p, arg1_offset, arg2_p,
+                                 arg2_offset, res_p, res_offset, depends);
 }
 
 template <typename fnT, typename T1, typename T2> struct NotEqualContigFactory
@@ -215,7 +193,7 @@ template <typename fnT, typename T1, typename T2> struct NotEqualContigFactory
 
 template <typename fnT, typename T1, typename T2> struct NotEqualTypeMapFactory
 {
-    /*! @brief get typeid for output type of operator()==(x, y), always bool */
+    /*! @brief get typeid for output type of operator()!=(x, y), always bool */
     std::enable_if_t<std::is_same<fnT, int>::value, int> get()
     {
         using rT = typename NotEqualOutputType<T1, T2>::value_type;
@@ -241,28 +219,11 @@ not_equal_strided_impl(sycl::queue exec_q,
                        const std::vector<sycl::event> &depends,
                        const std::vector<sycl::event> &additional_depends)
 {
-    sycl::event comp_ev = exec_q.submit([&](sycl::handler &cgh) {
-        cgh.depends_on(depends);
-        cgh.depends_on(additional_depends);
-
-        using resTy = typename NotEqualOutputType<argTy1, argTy2>::value_type;
-
-        using IndexerT =
-            typename dpctl::tensor::offset_utils::ThreeOffsets_StridedIndexer;
-
-        IndexerT indexer{nd, arg1_offset, arg2_offset, res_offset,
-                         shape_and_strides};
-
-        const argTy1 *arg1_tp = reinterpret_cast<const argTy1 *>(arg1_p);
-        const argTy2 *arg2_tp = reinterpret_cast<const argTy2 *>(arg2_p);
-        resTy *res_tp = reinterpret_cast<resTy *>(res_p);
-
-        cgh.parallel_for<
-            not_equal_strided_strided_kernel<argTy1, argTy2, resTy, IndexerT>>(
-            {nelems}, NotEqualStridedFunctor<argTy1, argTy2, resTy, IndexerT>(
-                          arg1_tp, arg2_tp, res_tp, indexer));
-    });
-    return comp_ev;
+    return elementwise_common::binary_strided_impl<
+        argTy1, argTy2, NotEqualOutputType, NotEqualStridedFunctor,
+        not_equal_strided_strided_kernel>(
+        exec_q, nelems, nd, shape_and_strides, arg1_p, arg1_offset, arg2_p,
+        arg2_offset, res_p, res_offset, depends, additional_depends);
 }
 
 template <typename fnT, typename T1, typename T2> struct NotEqualStridedFactory