Skip to content

Commit da64a4b

Browse files
antonwolfyvtavana
antonwolfy
and
vtavana
authored
Add implementation of dpnp.ldexp (#2110)
* Add _ldexp to ufunc extension * Add dpnp.ldexp to python interface * Enable third party test * Add CFD tests * Enable umath tests * Add support of `weak_type_resolver` keyword to DPNPBinaryFunc class * Add more integer type to matrix * Proper handling of large exponent values * Add more tests * Add a test with scalar exp of bool type * Fix exception raised in numpy.ldexp by test_overflow * Mute tests with unsupported loop for numpy.ldexp * Apply suggestions from code review Co-authored-by: vtavana <[email protected]> * State limitations for real, imag, rint, round, angle, fabs, fix functions * Remove limitation block for angle, real, round functions * Align real and imag signatures with ones from dpctl * Add missing description of deg keyword to dpnp.angle * State expected dtypes for input arrays * Fix typos in limitations --------- Co-authored-by: vtavana <[email protected]>
1 parent 89b03eb commit da64a4b

File tree

14 files changed

+526
-25
lines changed

14 files changed

+526
-25
lines changed

‎dpnp/backend/extensions/ufunc/CMakeLists.txt

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,7 @@ set(_elementwise_sources
3535
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/gcd.cpp
3636
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/heaviside.cpp
3737
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
38+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
3839
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
3940
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
4041
)

‎dpnp/backend/extensions/ufunc/elementwise_functions/common.cpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -35,6 +35,7 @@
3535
#include "gcd.hpp"
3636
#include "heaviside.hpp"
3737
#include "lcm.hpp"
38+
#include "ldexp.hpp"
3839
#include "logaddexp2.hpp"
3940
#include "radians.hpp"
4041

@@ -57,6 +58,7 @@ void init_elementwise_functions(py::module_ m)
5758
init_gcd(m);
5859
init_heaviside(m);
5960
init_lcm(m);
61+
init_ldexp(m);
6062
init_logaddexp2(m);
6163
init_radians(m);
6264
}

‎dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.cpp

Lines changed: 169 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,169 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// maxification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#include <sycl/sycl.hpp>
27+
28+
#include "dpctl4pybind11.hpp"
29+
30+
#include "kernels/elementwise_functions/ldexp.hpp"
31+
#include "ldexp.hpp"
32+
#include "populate.hpp"
33+
34+
// include a local copy of elementwise common header from dpctl tensor:
35+
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
36+
// TODO: replace by including dpctl header once available
37+
#include "../../elementwise_functions/elementwise_functions.hpp"
38+
39+
// dpctl tensor headers
40+
#include "kernels/elementwise_functions/common.hpp"
41+
#include "kernels/elementwise_functions/maximum.hpp"
42+
#include "utils/type_dispatch.hpp"
43+
44+
namespace dpnp::extensions::ufunc
45+
{
46+
namespace py = pybind11;
47+
namespace py_int = dpnp::extensions::py_internal;
48+
namespace td_ns = dpctl::tensor::type_dispatch;
49+
50+
namespace impl
51+
{
52+
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
53+
namespace max_ns = dpctl::tensor::kernels::maximum;
54+
55+
// Supports the same types table as for maximum function in dpctl
56+
// template <typename T1, typename T2>
57+
// using OutputType = max_ns::MaximumOutputType<T1, T2>;
58+
template <typename T1, typename T2>
59+
struct OutputType
60+
{
61+
using value_type = typename std::disjunction< // disjunction is C++17
62+
// feature, supported by DPC++
63+
td_ns::BinaryTypeMapResultEntry<T1,
64+
sycl::half,
65+
T2,
66+
std::int8_t,
67+
sycl::half>,
68+
td_ns::BinaryTypeMapResultEntry<T1,
69+
sycl::half,
70+
T2,
71+
std::int16_t,
72+
sycl::half>,
73+
td_ns::BinaryTypeMapResultEntry<T1,
74+
sycl::half,
75+
T2,
76+
std::int32_t,
77+
sycl::half>,
78+
td_ns::BinaryTypeMapResultEntry<T1,
79+
sycl::half,
80+
T2,
81+
std::int64_t,
82+
sycl::half>,
83+
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int8_t, float>,
84+
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int16_t, float>,
85+
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int32_t, float>,
86+
td_ns::BinaryTypeMapResultEntry<T1, float, T2, std::int64_t, float>,
87+
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int8_t, double>,
88+
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int16_t, double>,
89+
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int32_t, double>,
90+
td_ns::BinaryTypeMapResultEntry<T1, double, T2, std::int64_t, double>,
91+
td_ns::DefaultResultEntry<void>>::result_type;
92+
};
93+
94+
using dpnp::kernels::ldexp::LdexpFunctor;
95+
96+
template <typename argT1,
97+
typename argT2,
98+
typename resT,
99+
unsigned int vec_sz = 4,
100+
unsigned int n_vecs = 2,
101+
bool enable_sg_loadstore = true>
102+
using ContigFunctor =
103+
ew_cmn_ns::BinaryContigFunctor<argT1,
104+
argT2,
105+
resT,
106+
LdexpFunctor<argT1, argT2, resT>,
107+
vec_sz,
108+
n_vecs,
109+
enable_sg_loadstore>;
110+
111+
template <typename argT1, typename argT2, typename resT, typename IndexerT>
112+
using StridedFunctor =
113+
ew_cmn_ns::BinaryStridedFunctor<argT1,
114+
argT2,
115+
resT,
116+
IndexerT,
117+
LdexpFunctor<argT1, argT2, resT>>;
118+
119+
using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
120+
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
121+
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
122+
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
123+
124+
static binary_contig_impl_fn_ptr_t
125+
ldexp_contig_dispatch_table[td_ns::num_types][td_ns::num_types];
126+
static int ldexp_output_typeid_table[td_ns::num_types][td_ns::num_types];
127+
static binary_strided_impl_fn_ptr_t
128+
ldexp_strided_dispatch_table[td_ns::num_types][td_ns::num_types];
129+
130+
MACRO_POPULATE_DISPATCH_TABLES(ldexp);
131+
} // namespace impl
132+
133+
void init_ldexp(py::module_ m)
134+
{
135+
using arrayT = dpctl::tensor::usm_ndarray;
136+
using event_vecT = std::vector<sycl::event>;
137+
{
138+
impl::populate_ldexp_dispatch_tables();
139+
using impl::ldexp_contig_dispatch_table;
140+
using impl::ldexp_output_typeid_table;
141+
using impl::ldexp_strided_dispatch_table;
142+
143+
auto ldexp_pyapi = [&](const arrayT &src1, const arrayT &src2,
144+
const arrayT &dst, sycl::queue &exec_q,
145+
const event_vecT &depends = {}) {
146+
return py_int::py_binary_ufunc(
147+
src1, src2, dst, exec_q, depends, ldexp_output_typeid_table,
148+
ldexp_contig_dispatch_table, ldexp_strided_dispatch_table,
149+
// no support of C-contig row with broadcasting in OneMKL
150+
td_ns::NullPtrTable<
151+
impl::
152+
binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
153+
td_ns::NullPtrTable<
154+
impl::
155+
binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
156+
};
157+
m.def("_ldexp", ldexp_pyapi, "", py::arg("src1"), py::arg("src2"),
158+
py::arg("dst"), py::arg("sycl_queue"),
159+
py::arg("depends") = py::list());
160+
161+
auto ldexp_result_type_pyapi = [&](const py::dtype &dtype1,
162+
const py::dtype &dtype2) {
163+
return py_int::py_binary_ufunc_result_type(
164+
dtype1, dtype2, ldexp_output_typeid_table);
165+
};
166+
m.def("_ldexp_result_type", ldexp_result_type_pyapi);
167+
}
168+
}
169+
} // namespace dpnp::extensions::ufunc

‎dpnp/backend/extensions/ufunc/elementwise_functions/ldexp.hpp

Lines changed: 35 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,35 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#pragma once
27+
28+
#include <pybind11/pybind11.h>
29+
30+
namespace py = pybind11;
31+
32+
namespace dpnp::extensions::ufunc
33+
{
34+
void init_ldexp(py::module_ m);
35+
} // namespace dpnp::extensions::ufunc

‎dpnp/backend/kernels/elementwise_functions/ldexp.hpp

Lines changed: 55 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,55 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2024, Intel Corporation
3+
// All rights reserved.
4+
//
5+
// Redistribution and use in source and binary forms, with or without
6+
// modification, are permitted provided that the following conditions are met:
7+
// - Redistributions of source code must retain the above copyright notice,
8+
// this list of conditions and the following disclaimer.
9+
// - Redistributions in binary form must reproduce the above copyright notice,
10+
// this list of conditions and the following disclaimer in the documentation
11+
// and/or other materials provided with the distribution.
12+
//
13+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
14+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
15+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
16+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
17+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
18+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
19+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
20+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
21+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
22+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
23+
// THE POSSIBILITY OF SUCH DAMAGE.
24+
//*****************************************************************************
25+
26+
#pragma once
27+
28+
#include <sycl/sycl.hpp>
29+
30+
// dpctl tensor headers
31+
#include "utils/math_utils.hpp"
32+
#include "utils/type_utils.hpp"
33+
34+
namespace dpnp::kernels::ldexp
35+
{
36+
template <typename argT1, typename argT2, typename resT>
37+
struct LdexpFunctor
38+
{
39+
using supports_sg_loadstore = typename std::true_type;
40+
using supports_vec = typename std::false_type;
41+
42+
resT operator()(const argT1 &in1, const argT2 &in2) const
43+
{
44+
if (((int)in2) == in2) {
45+
return sycl::ldexp(in1, in2);
46+
}
47+
48+
// a separate handling for large integer values
49+
if (in2 > 0) {
50+
return std::numeric_limits<resT>::infinity();
51+
}
52+
return resT(0);
53+
}
54+
};
55+
} // namespace dpnp::kernels::ldexp

0 commit comments

Comments
 (0)