Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Rework implementation of dpnp.degrees and dpnp.rad2deg functions #1949

Merged
merged 4 commits into from
Jul 26, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
1 change: 1 addition & 0 deletions dpnp/backend/extensions/ufunc/CMakeLists.txt
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@

set(_elementwise_sources
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -25,6 +25,7 @@

#include <pybind11/pybind11.h>

#include "degrees.hpp"
#include "fabs.hpp"
#include "fmax.hpp"
#include "fmin.hpp"
Expand All @@ -40,6 +41,7 @@ namespace dpnp::extensions::ufunc
*/
void init_elementwise_functions(py::module_ m)
{
init_degrees(m);
init_fabs(m);
init_fmax(m);
init_fmin(m);
Expand Down
127 changes: 127 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/degrees.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,127 @@
//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#include <sycl/sycl.hpp>

#include "dpctl4pybind11.hpp"

#include "degrees.hpp"
#include "kernels/elementwise_functions/degrees.hpp"
#include "populate.hpp"

// include a local copy of elementwise common header from dpctl tensor:
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
// TODO: replace by including dpctl header once available
#include "../../elementwise_functions/elementwise_functions.hpp"

// dpctl tensor headers
#include "kernels/elementwise_functions/common.hpp"
#include "utils/type_dispatch.hpp"

namespace dpnp::extensions::ufunc
{
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;

namespace impl
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace td_ns = dpctl::tensor::type_dispatch;

/**
* @brief A factory to define pairs of supported types for which
* sycl::degrees<T> function is available.
*
* @tparam T Type of input vector `a` and of result vector `y`.
*/
template <typename T>
struct OutputType
{
using value_type =
typename std::disjunction<td_ns::TypeMapResultEntry<T, sycl::half>,
td_ns::TypeMapResultEntry<T, float>,
td_ns::TypeMapResultEntry<T, double>,
td_ns::DefaultResultEntry<void>>::result_type;
};

using dpnp::kernels::degrees::DegreesFunctor;

template <typename argT,
typename resT = argT,
unsigned int vec_sz = 4,
unsigned int n_vecs = 2,
bool enable_sg_loadstore = true>
using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
resT,
DegreesFunctor<argT, resT>,
vec_sz,
n_vecs,
enable_sg_loadstore>;

template <typename argTy, typename resTy, typename IndexerT>
using StridedFunctor = ew_cmn_ns::
UnaryStridedFunctor<argTy, resTy, IndexerT, DegreesFunctor<argTy, resTy>>;

using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;

static unary_contig_impl_fn_ptr_t
degrees_contig_dispatch_vector[td_ns::num_types];
static int degrees_output_typeid_vector[td_ns::num_types];
static unary_strided_impl_fn_ptr_t
degrees_strided_dispatch_vector[td_ns::num_types];

MACRO_POPULATE_DISPATCH_VECTORS(degrees);
} // namespace impl

void init_degrees(py::module_ m)
{
using arrayT = dpctl::tensor::usm_ndarray;
using event_vecT = std::vector<sycl::event>;
{
impl::populate_degrees_dispatch_vectors();
using impl::degrees_contig_dispatch_vector;
using impl::degrees_output_typeid_vector;
using impl::degrees_strided_dispatch_vector;

auto degrees_pyapi = [&](const arrayT &src, const arrayT &dst,
sycl::queue &exec_q,
const event_vecT &depends = {}) {
return py_int::py_unary_ufunc(src, dst, exec_q, depends,
degrees_output_typeid_vector,
degrees_contig_dispatch_vector,
degrees_strided_dispatch_vector);
};
m.def("_degrees", degrees_pyapi, "", py::arg("src"), py::arg("dst"),
py::arg("sycl_queue"), py::arg("depends") = py::list());

auto degrees_result_type_pyapi = [&](const py::dtype &dtype) {
return py_int::py_unary_ufunc_result_type(
dtype, degrees_output_typeid_vector);
};
m.def("_degrees_result_type", degrees_result_type_pyapi);
}
}
} // namespace dpnp::extensions::ufunc
35 changes: 35 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/degrees.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,35 @@
//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#pragma once

#include <pybind11/pybind11.h>

namespace py = pybind11;

namespace dpnp::extensions::ufunc
{
void init_degrees(py::module_ m);
} // namespace dpnp::extensions::ufunc
3 changes: 0 additions & 3 deletions dpnp/backend/include/dpnp_gen_1arg_2type_tbl.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -86,9 +86,6 @@
#endif

MACRO_1ARG_2TYPES_OP(dpnp_copyto_c, input_elem, q.submit(kernel_func))
MACRO_1ARG_2TYPES_OP(dpnp_degrees_c,
sycl::degrees(input_elem),
q.submit(kernel_func))
MACRO_1ARG_2TYPES_OP(dpnp_sqrt_c,
sycl::sqrt(input_elem),
oneapi::mkl::vm::sqrt(q, input1_size, input1_data, result))
Expand Down
3 changes: 0 additions & 3 deletions dpnp/backend/include/dpnp_iface_fptr.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -78,9 +78,6 @@ enum class DPNPFuncName : size_t
parameters */
DPNP_FN_COUNT_NONZERO, /**< Used in numpy.count_nonzero() impl */
DPNP_FN_COV, /**< Used in numpy.cov() impl */
DPNP_FN_DEGREES, /**< Used in numpy.degrees() impl */
DPNP_FN_DEGREES_EXT, /**< Used in numpy.degrees() impl, requires extra
parameters */
DPNP_FN_DOT, /**< Used in numpy.dot() impl */
DPNP_FN_DOT_EXT, /**< Used in numpy.dot() impl, requires extra parameters */
DPNP_FN_EDIFF1D, /**< Used in numpy.ediff1d() impl */
Expand Down
30 changes: 0 additions & 30 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -307,36 +307,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap)
fmap[DPNPFuncName::DPNP_FN_COPYTO_EXT][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_copyto_c_ext<float, float>};

fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_INT][eft_INT] = {
eft_DBL, (void *)dpnp_degrees_c_default<int32_t, double>};
fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_LNG][eft_LNG] = {
eft_DBL, (void *)dpnp_degrees_c_default<int64_t, double>};
fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_degrees_c_default<float, float>};
fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_degrees_c_default<double, double>};

fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_INT][eft_INT] = {
get_default_floating_type(),
(void *)dpnp_degrees_c_ext<
int32_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_degrees_c_ext<
int32_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_LNG][eft_LNG] = {
get_default_floating_type(),
(void *)dpnp_degrees_c_ext<
int64_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_degrees_c_ext<
int64_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_degrees_c_ext<float, float>};
fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_degrees_c_ext<double, double>};

fmap[DPNPFuncName::DPNP_FN_SQRT][eft_INT][eft_INT] = {
eft_DBL, (void *)dpnp_sqrt_c_default<int32_t, double>};
fmap[DPNPFuncName::DPNP_FN_SQRT][eft_LNG][eft_LNG] = {
Expand Down
55 changes: 55 additions & 0 deletions dpnp/backend/kernels/elementwise_functions/degrees.hpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,55 @@
//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
// this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
// this list of conditions and the following disclaimer in the documentation
// and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#pragma once

#include <sycl/sycl.hpp>

namespace dpnp::kernels::degrees
{
template <typename argT, typename resT>
struct DegreesFunctor
{
// 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::true_type;
// do both argT and resT support subgroup store/load operation
using supports_sg_loadstore = typename std::true_type;

resT operator()(const argT &x) const
{
return sycl::degrees(x);
}

template <int vec_sz>
sycl::vec<resT, vec_sz> operator()(const sycl::vec<argT, vec_sz> &x) const
{
return sycl::degrees(x);
}
};
} // namespace dpnp::kernels::degrees
7 changes: 0 additions & 7 deletions dpnp/dpnp_algo/dpnp_algo.pxd
Original file line number Diff line number Diff line change
Expand Up @@ -36,7 +36,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
DPNP_FN_ALLCLOSE_EXT
DPNP_FN_CHOOSE_EXT
DPNP_FN_CORRELATE_EXT
DPNP_FN_DEGREES_EXT
DPNP_FN_EDIFF1D_EXT
DPNP_FN_ERF_EXT
DPNP_FN_FFT_FFT_EXT
Expand Down Expand Up @@ -165,9 +164,3 @@ Logic functions
"""
cpdef dpnp_descriptor dpnp_isclose(dpnp_descriptor input1, dpnp_descriptor input2,
double rtol=*, double atol=*, cpp_bool equal_nan=*)


"""
Trigonometric functions
"""
cpdef dpnp_descriptor dpnp_degrees(dpnp_descriptor array1)
5 changes: 0 additions & 5 deletions dpnp/dpnp_algo/dpnp_algo_trigonometric.pxi
Original file line number Diff line number Diff line change
Expand Up @@ -36,15 +36,10 @@ and the rest of the library
# NO IMPORTs here. All imports must be placed into main "dpnp_algo.pyx" file

__all__ += [
'dpnp_degrees',
'dpnp_unwrap'
]


cpdef utils.dpnp_descriptor dpnp_degrees(utils.dpnp_descriptor x1):
return call_fptr_1in_1out_strides(DPNP_FN_DEGREES_EXT, x1)


cpdef utils.dpnp_descriptor dpnp_unwrap(utils.dpnp_descriptor array1):

result_type = dpnp.float64
Expand Down
Loading
Loading