Skip to content

Commit

Permalink
Rework implementation of dpnp.deg2rad and dpnp.radians functions (#…
Browse files Browse the repository at this point in the history
…1943)

* Add implementation of dpnp.radians() and dpnp.deg2rad()

* Add more tests

* Applied comments to fix typo in comment
antonwolfy authored Jul 26, 2024

Verified

This commit was signed with the committer’s verified signature.
Aaronontheweb Aaron Stannard
1 parent 8acbe64 commit e33a82b
Showing 21 changed files with 513 additions and 270 deletions.
4 changes: 4 additions & 0 deletions doc/reference/ufunc.rst
Original file line number Diff line number Diff line change
@@ -54,6 +54,8 @@ Math operations

Trigonometric functions
~~~~~~~~~~~~~~~~~~~~~~~
All trigonometric functions use radians when an angle is called for.
The ratio of degrees to radians is :math:`180^{\circ}/\pi.`

.. autosummary::
:toctree: generated/
@@ -73,6 +75,8 @@ Trigonometric functions
dpnp.arcsinh
dpnp.arccosh
dpnp.arctanh
dpnp.degrees
dpnp.radians
dpnp.deg2rad
dpnp.rad2deg

1 change: 1 addition & 0 deletions dpnp/backend/extensions/ufunc/CMakeLists.txt
Original file line number Diff line number Diff line change
@@ -29,6 +29,7 @@ set(_elementwise_sources
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
)

set(python_module_name _ufunc_impl)
Original file line number Diff line number Diff line change
@@ -29,6 +29,7 @@
#include "fmax.hpp"
#include "fmin.hpp"
#include "fmod.hpp"
#include "radians.hpp"

namespace py = pybind11;

@@ -43,5 +44,6 @@ void init_elementwise_functions(py::module_ m)
init_fmax(m);
init_fmin(m);
init_fmod(m);
init_radians(m);
}
} // namespace dpnp::extensions::ufunc
11 changes: 4 additions & 7 deletions dpnp/backend/extensions/ufunc/elementwise_functions/fabs.cpp
Original file line number Diff line number Diff line change
@@ -40,19 +40,16 @@
#include "kernels/elementwise_functions/common.hpp"
#include "utils/type_dispatch.hpp"

namespace py = pybind11;

namespace dpnp::extensions::ufunc
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;

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

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::fabs<T> function is available.
11 changes: 4 additions & 7 deletions dpnp/backend/extensions/ufunc/elementwise_functions/fmax.cpp
Original file line number Diff line number Diff line change
@@ -41,20 +41,17 @@
#include "kernels/elementwise_functions/maximum.hpp"
#include "utils/type_dispatch.hpp"

namespace py = pybind11;

namespace dpnp::extensions::ufunc
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace max_ns = dpctl::tensor::kernels::maximum;
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;

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

namespace impl
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace max_ns = dpctl::tensor::kernels::maximum;

// Supports the same types table as for maximum function in dpctl
template <typename T1, typename T2>
using OutputType = max_ns::MaximumOutputType<T1, T2>;
11 changes: 4 additions & 7 deletions dpnp/backend/extensions/ufunc/elementwise_functions/fmin.cpp
Original file line number Diff line number Diff line change
@@ -41,20 +41,17 @@
#include "kernels/elementwise_functions/minimum.hpp"
#include "utils/type_dispatch.hpp"

namespace py = pybind11;

namespace dpnp::extensions::ufunc
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace min_ns = dpctl::tensor::kernels::minimum;
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;

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

namespace impl
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace min_ns = dpctl::tensor::kernels::minimum;

// Supports the same types table as for minimum function in dpctl
template <typename T1, typename T2>
using OutputType = min_ns::MinimumOutputType<T1, T2>;
11 changes: 4 additions & 7 deletions dpnp/backend/extensions/ufunc/elementwise_functions/fmod.cpp
Original file line number Diff line number Diff line change
@@ -40,19 +40,16 @@
#include "kernels/elementwise_functions/common.hpp"
#include "utils/type_dispatch.hpp"

namespace py = pybind11;

namespace dpnp::extensions::ufunc
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;

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

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::fmod<T> function is available.
127 changes: 127 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/radians.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 "kernels/elementwise_functions/radians.hpp"
#include "populate.hpp"
#include "radians.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::radians<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::radians::RadiansFunctor;

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,
RadiansFunctor<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, RadiansFunctor<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
radians_contig_dispatch_vector[td_ns::num_types];
static int radians_output_typeid_vector[td_ns::num_types];
static unary_strided_impl_fn_ptr_t
radians_strided_dispatch_vector[td_ns::num_types];

MACRO_POPULATE_DISPATCH_VECTORS(radians);
} // namespace impl

void init_radians(py::module_ m)
{
using arrayT = dpctl::tensor::usm_ndarray;
using event_vecT = std::vector<sycl::event>;
{
impl::populate_radians_dispatch_vectors();
using impl::radians_contig_dispatch_vector;
using impl::radians_output_typeid_vector;
using impl::radians_strided_dispatch_vector;

auto radians_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,
radians_output_typeid_vector,
radians_contig_dispatch_vector,
radians_strided_dispatch_vector);
};
m.def("_radians", radians_pyapi, "", py::arg("src"), py::arg("dst"),
py::arg("sycl_queue"), py::arg("depends") = py::list());

auto radians_result_type_pyapi = [&](const py::dtype &dtype) {
return py_int::py_unary_ufunc_result_type(
dtype, radians_output_typeid_vector);
};
m.def("_radians_result_type", radians_result_type_pyapi);
}
}
} // namespace dpnp::extensions::ufunc
35 changes: 35 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/radians.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_radians(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
@@ -89,9 +89,6 @@ 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_radians_c,
sycl::radians(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))
3 changes: 0 additions & 3 deletions dpnp/backend/include/dpnp_iface_fptr.hpp
Original file line number Diff line number Diff line change
@@ -116,9 +116,6 @@ enum class DPNPFuncName : size_t
DPNP_FN_PARTITION_EXT, /**< Used in numpy.partition() impl, requires extra
parameters */
DPNP_FN_PROD, /**< Used in numpy.prod() impl */
DPNP_FN_RADIANS, /**< Used in numpy.radians() impl */
DPNP_FN_RADIANS_EXT, /**< Used in numpy.radians() impl, requires extra
parameters */
DPNP_FN_RNG_BETA, /**< Used in numpy.random.beta() impl */
DPNP_FN_RNG_BETA_EXT, /**< Used in numpy.random.beta() impl, requires extra
parameters */
30 changes: 0 additions & 30 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
Original file line number Diff line number Diff line change
@@ -337,36 +337,6 @@ static void func_map_init_elemwise_1arg_2type(func_map_t &fmap)
fmap[DPNPFuncName::DPNP_FN_DEGREES_EXT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_degrees_c_ext<double, double>};

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

fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_INT][eft_INT] = {
get_default_floating_type(),
(void *)dpnp_radians_c_ext<
int32_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_radians_c_ext<
int32_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_LNG][eft_LNG] = {
get_default_floating_type(),
(void *)dpnp_radians_c_ext<
int64_t, func_type_map_t::find_type<get_default_floating_type()>>,
get_default_floating_type<std::false_type>(),
(void *)dpnp_radians_c_ext<
int64_t, func_type_map_t::find_type<
get_default_floating_type<std::false_type>()>>};
fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_FLT][eft_FLT] = {
eft_FLT, (void *)dpnp_radians_c_ext<float, float>};
fmap[DPNPFuncName::DPNP_FN_RADIANS_EXT][eft_DBL][eft_DBL] = {
eft_DBL, (void *)dpnp_radians_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] = {
2 changes: 1 addition & 1 deletion dpnp/backend/kernels/elementwise_functions/fabs.hpp
Original file line number Diff line number Diff line change
@@ -38,7 +38,7 @@ struct FabsFunctor
// constexpr resT constant_value = resT{};
// is function defined for sycl::vec
using supports_vec = typename std::false_type;
// do both argT and resT support sugroup store/load operation
// do both argT and resT support subgroup store/load operation
using supports_sg_loadstore = typename std::true_type;

resT operator()(const argT &x) const
Loading

0 comments on commit e33a82b

Please sign in to comment.