Skip to content

Add implementation of dpnp.ldexp #2110

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

Merged
merged 23 commits into from
Oct 18, 2024
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
23 commits
Select commit Hold shift + click to select a range
ee89a6e
Add _ldexp to ufunc extension
antonwolfy Oct 9, 2024
12ccad5
Add dpnp.ldexp to python interface
antonwolfy Oct 9, 2024
0fcba0a
Enable third party test
antonwolfy Oct 9, 2024
219c243
Add CFD tests
antonwolfy Oct 9, 2024
0bb6761
Enable umath tests
antonwolfy Oct 9, 2024
467e44b
Add support of `weak_type_resolver` keyword to DPNPBinaryFunc class
antonwolfy Oct 9, 2024
c3bb14f
Add more integer type to matrix
antonwolfy Oct 14, 2024
65a75fb
Proper handling of large exponent values
antonwolfy Oct 14, 2024
a107236
Add more tests
antonwolfy Oct 14, 2024
afb20b0
Add a test with scalar exp of bool type
antonwolfy Oct 14, 2024
4acb885
Fix exception raised in numpy.ldexp by test_overflow
antonwolfy Oct 15, 2024
ce507b1
Merge branch 'master' into impl-ldexp
antonwolfy Oct 15, 2024
9fcca48
Mute tests with unsupported loop for numpy.ldexp
antonwolfy Oct 15, 2024
710c2e7
Merge branch 'master' into impl-ldexp
antonwolfy Oct 15, 2024
7cabd1b
Merge branch 'master' into impl-ldexp
antonwolfy Oct 15, 2024
b6ba66e
Apply suggestions from code review
antonwolfy Oct 17, 2024
0d91c94
State limitations for real, imag, rint, round, angle, fabs, fix funct…
antonwolfy Oct 17, 2024
ac69f42
Remove limitation block for angle, real, round functions
antonwolfy Oct 17, 2024
8bdfd40
Align real and imag signatures with ones from dpctl
antonwolfy Oct 17, 2024
54aab90
Add missing description of deg keyword to dpnp.angle
antonwolfy Oct 17, 2024
14f564c
Merge branch 'master' into impl-ldexp
antonwolfy Oct 18, 2024
5a35f78
State expected dtypes for input arrays
antonwolfy Oct 18, 2024
7414487
Fix typos in limitations
antonwolfy Oct 18, 2024
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 @@ -35,6 +35,7 @@ set(_elementwise_sources
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/gcd.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
)
Expand Down
Original file line number Diff line number Diff line change
Expand Up @@ -35,6 +35,7 @@
#include "gcd.hpp"
#include "heaviside.hpp"
#include "lcm.hpp"
#include "ldexp.hpp"
#include "logaddexp2.hpp"
#include "radians.hpp"

Expand All @@ -57,6 +58,7 @@ void init_elementwise_functions(py::module_ m)
init_gcd(m);
init_heaviside(m);
init_lcm(m);
init_ldexp(m);
init_logaddexp2(m);
init_radians(m);
}
Expand Down
169 changes: 169 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp
Original file line number Diff line number Diff line change
@@ -0,0 +1,169 @@
//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// maxification, 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/ldexp.hpp"
#include "ldexp.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 "kernels/elementwise_functions/maximum.hpp"
#include "utils/type_dispatch.hpp"

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

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>;
template <typename T1, typename T2>
struct OutputType
{
using value_type = typename std::disjunction< // disjunction is C++17
// feature, supported by DPC++
td_ns::BinaryTypeMapResultEntry<T1,
sycl::half,
T2,
std::int8_t,
sycl::half>,
td_ns::BinaryTypeMapResultEntry<T1,
sycl::half,
T2,
std::int16_t,
sycl::half>,
td_ns::BinaryTypeMapResultEntry<T1,
sycl::half,
T2,
std::int32_t,
sycl::half>,
td_ns::BinaryTypeMapResultEntry<T1,
sycl::half,
T2,
std::int64_t,
sycl::half>,
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int8_t, float>,
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int16_t, float>,
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int32_t, float>,
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int64_t, float>,
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int8_t, double>,
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int16_t, double>,
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int32_t, double>,
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int64_t, double>,
td_ns::DefaultResultEntry<void>>::result_type;
};

using dpnp::kernels::ldexp::LdexpFunctor;

template <typename argT1,
typename argT2,
typename resT,
unsigned int vec_sz = 4,
unsigned int n_vecs = 2,
bool enable_sg_loadstore = true>
using ContigFunctor =
ew_cmn_ns::BinaryContigFunctor<argT1,
argT2,
resT,
LdexpFunctor<argT1, argT2, resT>,
vec_sz,
n_vecs,
enable_sg_loadstore>;

template <typename argT1, typename argT2, typename resT, typename IndexerT>
using StridedFunctor =
ew_cmn_ns::BinaryStridedFunctor<argT1,
argT2,
resT,
IndexerT,
LdexpFunctor<argT1, argT2, resT>>;

using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;

static binary_contig_impl_fn_ptr_t
ldexp_contig_dispatch_table[td_ns::num_types][td_ns::num_types];
static int ldexp_output_typeid_table[td_ns::num_types][td_ns::num_types];
static binary_strided_impl_fn_ptr_t
ldexp_strided_dispatch_table[td_ns::num_types][td_ns::num_types];

MACRO_POPULATE_DISPATCH_TABLES(ldexp);
} // namespace impl

void init_ldexp(py::module_ m)
{
using arrayT = dpctl::tensor::usm_ndarray;
using event_vecT = std::vector<sycl::event>;
{
impl::populate_ldexp_dispatch_tables();
using impl::ldexp_contig_dispatch_table;
using impl::ldexp_output_typeid_table;
using impl::ldexp_strided_dispatch_table;

auto ldexp_pyapi = [&](const arrayT &src1, const arrayT &src2,
const arrayT &dst, sycl::queue &exec_q,
const event_vecT &depends = {}) {
return py_int::py_binary_ufunc(
src1, src2, dst, exec_q, depends, ldexp_output_typeid_table,
ldexp_contig_dispatch_table, ldexp_strided_dispatch_table,
// no support of C-contig row with broadcasting in OneMKL
td_ns::NullPtrTable<
impl::
binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
td_ns::NullPtrTable<
impl::
binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
};
m.def("_ldexp", ldexp_pyapi, "", py::arg("src1"), py::arg("src2"),
py::arg("dst"), py::arg("sycl_queue"),
py::arg("depends") = py::list());

auto ldexp_result_type_pyapi = [&](const py::dtype &dtype1,
const py::dtype &dtype2) {
return py_int::py_binary_ufunc_result_type(
dtype1, dtype2, ldexp_output_typeid_table);
};
m.def("_ldexp_result_type", ldexp_result_type_pyapi);
}
}
} // namespace dpnp::extensions::ufunc
35 changes: 35 additions & 0 deletions dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.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_ldexp(py::module_ m);
} // namespace dpnp::extensions::ufunc
55 changes: 55 additions & 0 deletions dpnp/backend/kernels/elementwise_functions/ldexp.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>

// dpctl tensor headers
#include "utils/math_utils.hpp"
#include "utils/type_utils.hpp"

namespace dpnp::kernels::ldexp
{
template <typename argT1, typename argT2, typename resT>
struct LdexpFunctor
{
using supports_sg_loadstore = typename std::true_type;
using supports_vec = typename std::false_type;

resT operator()(const argT1 &in1, const argT2 &in2) const
{
if (((int)in2) == in2) {
return sycl::ldexp(in1, in2);
}

// a separate handling for large integer values
if (in2 > 0) {
return std::numeric_limits<resT>::infinity();
}
return resT(0);
}
};
} // namespace dpnp::kernels::ldexp
Loading
Loading