Skip to content

Commit 4082593

Browse files
antonwolfyantonwolfy
authored
Redesigned dpnp.modf function to be a part of ufunc and vm pybind11 extensions (#2654)
The PR updates implementation of `dpnp.modf` to move it from old legacy backend to be part of `ufunc` and `vm` extensions. Also `DPNPUnaryTwoOutputsFunc` class is extended with support of oneMKL VM functions. The tests for unary functions with 2 output arrays are separated to `dpnp/tests/test_unary_two_outputs_ufuncs.py` file. And the legacy implementation of `dpnp.modf` was cleaned up.
1 parent 6f6855b commit 4082593

35 files changed

+994
-648
lines changed

‎CHANGELOG.md‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -31,6 +31,7 @@ Also, that release drops support for Python 3.9, making Python 3.10 the minimum
3131
* Defined explicit versions range of the Python interpreter which is needed during the build [#2634](https://github.com/IntelPython/dpnp/pull/2634)
3232
* Aligned documentation with NumPy and CuPy style by using short function names [#2633](https://github.com/IntelPython/dpnp/pull/2633)
3333
* Added the missing positional-only and keyword-only parameter markers to bring the ufunc signatures into alignment with NumPy [#2660](https://github.com/IntelPython/dpnp/pull/2660)
34+
* Redesigned `dpnp.modf` function to be a part of `ufunc` and `vm` pybind11 extensions [#2654](https://github.com/IntelPython/dpnp/pull/2654)
3435

3536
### Deprecated
3637

‎dpnp/backend/CMakeLists.txt‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -30,7 +30,6 @@
3030
set(DPNP_SRC
3131
kernels/dpnp_krnl_arraycreation.cpp
3232
kernels/dpnp_krnl_common.cpp
33-
kernels/dpnp_krnl_mathematical.cpp
3433
kernels/dpnp_krnl_random.cpp
3534
kernels/dpnp_krnl_sorting.cpp
3635
src/dpnp_iface_fptr.cpp

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -47,6 +47,7 @@ set(_elementwise_sources
4747
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/lcm.cpp
4848
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/ldexp.cpp
4949
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/logaddexp2.cpp
50+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/modf.cpp
5051
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/nan_to_num.cpp
5152
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/radians.cpp
5253
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/sinc.cpp

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -46,6 +46,7 @@
4646
#include "lcm.hpp"
4747
#include "ldexp.hpp"
4848
#include "logaddexp2.hpp"
49+
#include "modf.hpp"
4950
#include "nan_to_num.hpp"
5051
#include "radians.hpp"
5152
#include "sinc.hpp"
@@ -78,6 +79,7 @@ void init_elementwise_functions(py::module_ m)
7879
init_lcm(m);
7980
init_ldexp(m);
8081
init_logaddexp2(m);
82+
init_modf(m);
8183
init_nan_to_num(m);
8284
init_radians(m);
8385
init_sinc(m);

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

Lines changed: 5 additions & 116 deletions
Original file line numberDiff line numberDiff line change
@@ -37,6 +37,7 @@
3737

3838
#include "frexp.hpp"
3939
#include "kernels/elementwise_functions/frexp.hpp"
40+
#include "populate.hpp"
4041

4142
// include a local copy of elementwise common header from dpctl tensor:
4243
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
@@ -65,10 +66,9 @@ namespace td_int_ns = py_int::type_dispatch;
6566
namespace td_ns = dpctl::tensor::type_dispatch;
6667

6768
using dpnp::kernels::frexp::FrexpFunctor;
68-
using ext::common::init_dispatch_vector;
6969

7070
template <typename T>
71-
struct FrexpOutputType
71+
struct OutputType
7272
{
7373
using table_type = std::disjunction< // disjunction is C++17
7474
// feature, supported by DPC++
@@ -81,15 +81,13 @@ struct FrexpOutputType
8181
using value_type2 = typename table_type::result_type2;
8282
};
8383

84-
// contiguous implementation
85-
8684
template <typename argTy,
8785
typename resTy1 = argTy,
8886
typename resTy2 = argTy,
8987
std::uint8_t vec_sz = 4u,
9088
std::uint8_t n_vecs = 2u,
9189
bool enable_sg_loadstore = true>
92-
using FrexpContigFunctor =
90+
using ContigFunctor =
9391
ew_cmn_ns::UnaryTwoOutputsContigFunctor<argTy,
9492
resTy1,
9593
resTy2,
@@ -98,115 +96,14 @@ using FrexpContigFunctor =
9896
n_vecs,
9997
enable_sg_loadstore>;
10098

101-
// strided implementation
102-
10399
template <typename argTy, typename resTy1, typename resTy2, typename IndexerT>
104-
using FrexpStridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
100+
using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
105101
argTy,
106102
resTy1,
107103
resTy2,
108104
IndexerT,
109105
FrexpFunctor<argTy, resTy1, resTy2>>;
110106

111-
template <typename T1,
112-
typename T2,
113-
typename T3,
114-
unsigned int vec_sz,
115-
unsigned int n_vecs>
116-
class frexp_contig_kernel;
117-
118-
template <typename argTy>
119-
sycl::event frexp_contig_impl(sycl::queue &exec_q,
120-
size_t nelems,
121-
const char *arg_p,
122-
char *res1_p,
123-
char *res2_p,
124-
const std::vector<sycl::event> &depends = {})
125-
{
126-
return ew_cmn_ns::unary_two_outputs_contig_impl<
127-
argTy, FrexpOutputType, FrexpContigFunctor, frexp_contig_kernel>(
128-
exec_q, nelems, arg_p, res1_p, res2_p, depends);
129-
}
130-
131-
template <typename fnT, typename T>
132-
struct FrexpContigFactory
133-
{
134-
fnT get()
135-
{
136-
if constexpr (std::is_same_v<typename FrexpOutputType<T>::value_type1,
137-
void> ||
138-
std::is_same_v<typename FrexpOutputType<T>::value_type2,
139-
void>)
140-
{
141-
fnT fn = nullptr;
142-
return fn;
143-
}
144-
else {
145-
fnT fn = frexp_contig_impl<T>;
146-
return fn;
147-
}
148-
}
149-
};
150-
151-
template <typename T1, typename T2, typename T3, typename T4>
152-
class frexp_strided_kernel;
153-
154-
template <typename argTy>
155-
sycl::event
156-
frexp_strided_impl(sycl::queue &exec_q,
157-
size_t nelems,
158-
int nd,
159-
const ssize_t *shape_and_strides,
160-
const char *arg_p,
161-
ssize_t arg_offset,
162-
char *res1_p,
163-
ssize_t res1_offset,
164-
char *res2_p,
165-
ssize_t res2_offset,
166-
const std::vector<sycl::event> &depends,
167-
const std::vector<sycl::event> &additional_depends)
168-
{
169-
return ew_cmn_ns::unary_two_outputs_strided_impl<
170-
argTy, FrexpOutputType, FrexpStridedFunctor, frexp_strided_kernel>(
171-
exec_q, nelems, nd, shape_and_strides, arg_p, arg_offset, res1_p,
172-
res1_offset, res2_p, res2_offset, depends, additional_depends);
173-
}
174-
175-
template <typename fnT, typename T>
176-
struct FrexpStridedFactory
177-
{
178-
fnT get()
179-
{
180-
if constexpr (std::is_same_v<typename FrexpOutputType<T>::value_type1,
181-
void> ||
182-
std::is_same_v<typename FrexpOutputType<T>::value_type2,
183-
void>)
184-
{
185-
fnT fn = nullptr;
186-
return fn;
187-
}
188-
else {
189-
fnT fn = frexp_strided_impl<T>;
190-
return fn;
191-
}
192-
}
193-
};
194-
195-
template <typename fnT, typename T>
196-
struct FrexpTypeMapFactory
197-
{
198-
/*! @brief get typeid for output type of sycl::frexp(T x) */
199-
std::enable_if_t<std::is_same<fnT, std::pair<int, int>>::value,
200-
std::pair<int, int>>
201-
get()
202-
{
203-
using rT1 = typename FrexpOutputType<T>::value_type1;
204-
using rT2 = typename FrexpOutputType<T>::value_type2;
205-
return std::make_pair(td_ns::GetTypeid<rT1>{}.get(),
206-
td_ns::GetTypeid<rT2>{}.get());
207-
}
208-
};
209-
210107
using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t;
211108
using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t;
212109

@@ -216,15 +113,7 @@ static std::pair<int, int> frexp_output_typeid_vector[td_ns::num_types];
216113
static unary_two_outputs_strided_impl_fn_ptr_t
217114
frexp_strided_dispatch_vector[td_ns::num_types];
218115

219-
void populate_frexp_dispatch_vectors(void)
220-
{
221-
init_dispatch_vector<unary_two_outputs_contig_impl_fn_ptr_t,
222-
FrexpContigFactory>(frexp_contig_dispatch_vector);
223-
init_dispatch_vector<unary_two_outputs_strided_impl_fn_ptr_t,
224-
FrexpStridedFactory>(frexp_strided_dispatch_vector);
225-
init_dispatch_vector<std::pair<int, int>, FrexpTypeMapFactory>(
226-
frexp_output_typeid_vector);
227-
};
116+
MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(frexp);
228117
} // namespace impl
229118

230119
void init_frexp(py::module_ m)

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

Lines changed: 146 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,146 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2025, 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+
// - Neither the name of the copyright holder nor the names of its contributors
13+
// may be used to endorse or promote products derived from this software
14+
// without specific prior written permission.
15+
//
16+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
17+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
19+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
20+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
21+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
22+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
23+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
24+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
25+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
26+
// THE POSSIBILITY OF SUCH DAMAGE.
27+
//*****************************************************************************
28+
29+
#include <cstdint>
30+
#include <type_traits>
31+
#include <utility>
32+
#include <vector>
33+
34+
#include <sycl/sycl.hpp>
35+
36+
#include "dpctl4pybind11.hpp"
37+
38+
#include "kernels/elementwise_functions/modf.hpp"
39+
#include "modf.hpp"
40+
#include "populate.hpp"
41+
42+
// include a local copy of elementwise common header from dpctl tensor:
43+
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
44+
// TODO: replace by including dpctl header once available
45+
#include "../../elementwise_functions/elementwise_functions.hpp"
46+
47+
#include "../../elementwise_functions/common.hpp"
48+
#include "../../elementwise_functions/type_dispatch_building.hpp"
49+
50+
// utils extension header
51+
#include "ext/common.hpp"
52+
53+
// dpctl tensor headers
54+
#include "kernels/elementwise_functions/common.hpp"
55+
#include "utils/type_dispatch.hpp"
56+
57+
namespace dpnp::extensions::ufunc
58+
{
59+
namespace py = pybind11;
60+
namespace py_int = dpnp::extensions::py_internal;
61+
62+
namespace impl
63+
{
64+
namespace ew_cmn_ns = dpnp::extensions::py_internal::elementwise_common;
65+
namespace td_int_ns = py_int::type_dispatch;
66+
namespace td_ns = dpctl::tensor::type_dispatch;
67+
68+
using dpnp::kernels::modf::ModfFunctor;
69+
70+
template <typename T>
71+
struct OutputType
72+
{
73+
using table_type = std::disjunction< // disjunction is C++17
74+
// feature, supported by DPC++
75+
td_int_ns::TypeMapTwoResultsEntry<T, sycl::half>,
76+
td_int_ns::TypeMapTwoResultsEntry<T, float>,
77+
td_int_ns::TypeMapTwoResultsEntry<T, double>,
78+
td_int_ns::DefaultTwoResultsEntry<void>>;
79+
using value_type1 = typename table_type::result_type1;
80+
using value_type2 = typename table_type::result_type2;
81+
};
82+
83+
template <typename argTy,
84+
typename resTy1 = argTy,
85+
typename resTy2 = argTy,
86+
std::uint8_t vec_sz = 4u,
87+
std::uint8_t n_vecs = 2u,
88+
bool enable_sg_loadstore = true>
89+
using ContigFunctor =
90+
ew_cmn_ns::UnaryTwoOutputsContigFunctor<argTy,
91+
resTy1,
92+
resTy2,
93+
ModfFunctor<argTy, resTy1, resTy2>,
94+
vec_sz,
95+
n_vecs,
96+
enable_sg_loadstore>;
97+
98+
template <typename argTy, typename resTy1, typename resTy2, typename IndexerT>
99+
using StridedFunctor = ew_cmn_ns::UnaryTwoOutputsStridedFunctor<
100+
argTy,
101+
resTy1,
102+
resTy2,
103+
IndexerT,
104+
ModfFunctor<argTy, resTy1, resTy2>>;
105+
106+
using ew_cmn_ns::unary_two_outputs_contig_impl_fn_ptr_t;
107+
using ew_cmn_ns::unary_two_outputs_strided_impl_fn_ptr_t;
108+
109+
static unary_two_outputs_contig_impl_fn_ptr_t
110+
modf_contig_dispatch_vector[td_ns::num_types];
111+
static std::pair<int, int> modf_output_typeid_vector[td_ns::num_types];
112+
static unary_two_outputs_strided_impl_fn_ptr_t
113+
modf_strided_dispatch_vector[td_ns::num_types];
114+
115+
MACRO_POPULATE_DISPATCH_2OUTS_VECTORS(modf);
116+
} // namespace impl
117+
118+
void init_modf(py::module_ m)
119+
{
120+
using arrayT = dpctl::tensor::usm_ndarray;
121+
using event_vecT = std::vector<sycl::event>;
122+
{
123+
impl::populate_modf_dispatch_vectors();
124+
using impl::modf_contig_dispatch_vector;
125+
using impl::modf_output_typeid_vector;
126+
using impl::modf_strided_dispatch_vector;
127+
128+
auto modf_pyapi = [&](const arrayT &src, const arrayT &dst1,
129+
const arrayT &dst2, sycl::queue &exec_q,
130+
const event_vecT &depends = {}) {
131+
return py_int::py_unary_two_outputs_ufunc(
132+
src, dst1, dst2, exec_q, depends, modf_output_typeid_vector,
133+
modf_contig_dispatch_vector, modf_strided_dispatch_vector);
134+
};
135+
m.def("_modf", modf_pyapi, "", py::arg("src"), py::arg("dst1"),
136+
py::arg("dst2"), py::arg("sycl_queue"),
137+
py::arg("depends") = py::list());
138+
139+
auto modf_result_type_pyapi = [&](const py::dtype &dtype) {
140+
return py_int::py_unary_two_outputs_ufunc_result_type(
141+
dtype, modf_output_typeid_vector);
142+
};
143+
m.def("_modf_result_type", modf_result_type_pyapi);
144+
}
145+
}
146+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2025, 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+
// - Neither the name of the copyright holder nor the names of its contributors
13+
// may be used to endorse or promote products derived from this software
14+
// without specific prior written permission.
15+
//
16+
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
17+
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18+
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
19+
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
20+
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
21+
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
22+
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
23+
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
24+
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
25+
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
26+
// THE POSSIBILITY OF SUCH DAMAGE.
27+
//*****************************************************************************
28+
29+
#pragma once
30+
31+
#include <pybind11/pybind11.h>
32+
33+
namespace py = pybind11;
34+
35+
namespace dpnp::extensions::ufunc
36+
{
37+
void init_modf(py::module_ m);
38+
} // namespace dpnp::extensions::ufunc

0 commit comments

Comments
 (0)