impl_real_imag_conj

Author: vtavana

dpctl/tensor/__init__.py

@@ -94,14 +94,17 @@
 from ._elementwise_funcs import (
     abs,
     add,
+    conj,
     cos,
     divide,
     equal,
     exp,
+    imag,
     isfinite,
     isinf,
     isnan,
     multiply,
+    real,
     sin,
     sqrt,
     subtract,
@@ -195,4 +198,7 @@
     "multiply",
     "subtract",
     "equal",
+    "real",
+    "imag",
+    "conj",
 ]

dpctl/tensor/_elementwise_funcs.py

@@ -90,7 +90,15 @@
 # FIXME: implement U09
 
 # U10: ==== CONJ          (x)
-# FIXME: implement U10
+_conj_docstring = """
+conj(x, order='K')
+
+Computes conjugate of each element `x_i` for input array `x`.
+"""
+
+conj = UnaryElementwiseFunc(
+    "conj", ti._conj_result_type, ti._conj, _conj_docstring
+)
 
 # U11: ==== COS           (x)
 _cos_docstring = """
@@ -174,7 +182,15 @@
 # FIXME: implement B12
 
 # U16: ==== IMAG        (x)
-# FIXME: implement U16
+_imag_docstring = """
+imag(x, order='K')
+
+Computes imaginary part of each element `x_i` for input array `x`.
+"""
+
+imag = UnaryElementwiseFunc(
+    "imag", ti._imag_result_type, ti._imag, _imag_docstring
+)
 
 # U17: ==== ISFINITE    (x)
 _isfinite_docstring_ = """
@@ -276,7 +292,15 @@
 # FIXME: implement B21
 
 # U27: ==== REAL        (x)
-# FIXME: implement U27
+_real_docstring = """
+real(x, order='K')
+
+Computes real part of each element `x_i` for input array `x`.
+"""
+
+real = UnaryElementwiseFunc(
+    "real", ti._real_result_type, ti._real, _real_docstring
+)
 
 # B22: ==== REMAINDER   (x1, x2)
 # FIXME: implement B22

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

@@ -0,0 +1,194 @@
+//=== conj.hpp -   Unary function CONJ                     ------
+//*-C++-*--/===//
+//
+//                      Data Parallel Control (dpctl)
+//
+// Copyright 2020-2023 Intel Corporation
+//
+// Licensed under the Apache License, Version 2.0 (the "License");
+// you may not use this file except in compliance with the License.
+// You may obtain a copy of the License at
+//
+//    http://www.apache.org/licenses/LICENSE-2.0
+//
+// Unless required by applicable law or agreed to in writing, software
+// distributed under the License is distributed on an "AS IS" BASIS,
+// WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+// See the License for the specific language governing permissions and
+// limitations under the License.
+//
+//===---------------------------------------------------------------------===//
+///
+/// \file
+/// This file defines kernels for elementwise evaluation of CONJ(x) function.
+//===---------------------------------------------------------------------===//
+
+#pragma once
+#include <CL/sycl.hpp>
+#include <cmath>
+#include <complex>
+#include <cstddef>
+#include <cstdint>
+#include <type_traits>
+
+#include "kernels/elementwise_functions/common.hpp"
+
+#include "utils/offset_utils.hpp"
+#include "utils/type_dispatch.hpp"
+#include "utils/type_utils.hpp"
+#include <pybind11/pybind11.h>
+
+namespace dpctl
+{
+namespace tensor
+{
+namespace kernels
+{
+namespace conj
+{
+
+namespace py = pybind11;
+namespace td_ns = dpctl::tensor::type_dispatch;
+
+using dpctl::tensor::type_utils::is_complex;
+
+template <typename argT, typename resT> struct ConjFunctor
+{
+
+    // is function constant for given argT
+    using is_constant = typename std::false_type;
+    // constant value, if constant
+    // constexpr resT constant_value = resT{};
+    // is function defined for sycl::vec
+    using supports_vec = typename std::false_type;
+    // do both argTy and resTy support sugroup store/load operation
+    using supports_sg_loadstore = typename std::negation<
+        std::disjunction<is_complex<resT>, is_complex<argT>>>;
+
+    resT operator()(const argT &in)
+    {
+        if constexpr (is_complex<argT>::value) {
+            return std::conj(in);
+        }
+        else {
+            if constexpr (!std::is_same_v<argT, bool>)
+                static_assert(std::is_same_v<resT, argT>);
+            return in;
+        }
+    }
+};
+
+template <typename argTy,
+          typename resTy = argTy,
+          unsigned int vec_sz = 4,
+          unsigned int n_vecs = 2>
+using ConjContigFunctor = elementwise_common::
+    UnaryContigFunctor<argTy, resTy, ConjFunctor<argTy, resTy>, vec_sz, n_vecs>;
+
+template <typename argTy, typename resTy, typename IndexerT>
+using ConjStridedFunctor = elementwise_common::
+    UnaryStridedFunctor<argTy, resTy, IndexerT, ConjFunctor<argTy, resTy>>;
+
+template <typename T> struct ConjOutputType
+{
+    using value_type = typename std::disjunction< // disjunction is C++17
+                                                  // feature, supported by DPC++
+        td_ns::TypeMapResultEntry<T, bool, int8_t>,
+        td_ns::TypeMapResultEntry<T, std::uint8_t>,
+        td_ns::TypeMapResultEntry<T, std::uint16_t>,
+        td_ns::TypeMapResultEntry<T, std::uint32_t>,
+        td_ns::TypeMapResultEntry<T, std::uint64_t>,
+        td_ns::TypeMapResultEntry<T, std::int8_t>,
+        td_ns::TypeMapResultEntry<T, std::int16_t>,
+        td_ns::TypeMapResultEntry<T, std::int32_t>,
+        td_ns::TypeMapResultEntry<T, std::int64_t>,
+        td_ns::TypeMapResultEntry<T, sycl::half>,
+        td_ns::TypeMapResultEntry<T, float>,
+        td_ns::TypeMapResultEntry<T, double>,
+        td_ns::TypeMapResultEntry<T, std::complex<float>>,
+        td_ns::TypeMapResultEntry<T, std::complex<double>>,
+        td_ns::DefaultResultEntry<void>>::result_type;
+};
+
+template <typename T1, typename T2, unsigned int vec_sz, unsigned int n_vecs>
+class conj_contig_kernel;
+
+template <typename argTy>
+sycl::event conj_contig_impl(sycl::queue exec_q,
+                             size_t nelems,
+                             const char *arg_p,
+                             char *res_p,
+                             const std::vector<sycl::event> &depends = {})
+{
+    return elementwise_common::unary_contig_impl<
+        argTy, ConjOutputType, ConjContigFunctor, conj_contig_kernel>(
+        exec_q, nelems, arg_p, res_p, depends);
+}
+
+template <typename fnT, typename T> struct ConjContigFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<typename ConjOutputType<T>::value_type,
+                                     void>) {
+            fnT fn = nullptr;
+            return fn;
+        }
+        else {
+            fnT fn = conj_contig_impl<T>;
+            return fn;
+        }
+    }
+};
+
+template <typename fnT, typename T> struct ConjTypeMapFactory
+{
+    /*! @brief get typeid for output type of std::conj(T x) */
+    std::enable_if_t<std::is_same<fnT, int>::value, int> get()
+    {
+        using rT = typename ConjOutputType<T>::value_type;
+        return td_ns::GetTypeid<rT>{}.get();
+    }
+};
+
+template <typename T1, typename T2, typename T3> class conj_strided_kernel;
+
+template <typename argTy>
+sycl::event
+conj_strided_impl(sycl::queue exec_q,
+                  size_t nelems,
+                  int nd,
+                  const py::ssize_t *shape_and_strides,
+                  const char *arg_p,
+                  py::ssize_t arg_offset,
+                  char *res_p,
+                  py::ssize_t res_offset,
+                  const std::vector<sycl::event> &depends,
+                  const std::vector<sycl::event> &additional_depends)
+{
+    return elementwise_common::unary_strided_impl<
+        argTy, ConjOutputType, ConjStridedFunctor, conj_strided_kernel>(
+        exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res_p,
+        res_offset, depends, additional_depends);
+}
+
+template <typename fnT, typename T> struct ConjStridedFactory
+{
+    fnT get()
+    {
+        if constexpr (std::is_same_v<typename ConjOutputType<T>::value_type,
+                                     void>) {
+            fnT fn = nullptr;
+            return fn;
+        }
+        else {
+            fnT fn = conj_strided_impl<T>;
+            return fn;
+        }
+    }
+};
+
+} // namespace conj
+} // namespace kernels
+} // namespace tensor
+} // namespace dpctl