Skip to content

Commit 0e10a07

Browse files
antonwolfyantonwolfy
committed
Implement dpnp.fix()
1 parent 543605c commit 0e10a07

File tree

6 files changed

+280
-0
lines changed

6 files changed

+280
-0
lines changed

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

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@ set(_elementwise_sources
2727
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/common.cpp
2828
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/degrees.cpp
2929
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fabs.cpp
30+
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fix.cpp
3031
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
3132
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.cpp
3233
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmod.cpp

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

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -27,6 +27,7 @@
2727

2828
#include "degrees.hpp"
2929
#include "fabs.hpp"
30+
#include "fix.hpp"
3031
#include "fmax.hpp"
3132
#include "fmin.hpp"
3233
#include "fmod.hpp"
@@ -44,6 +45,7 @@ void init_elementwise_functions(py::module_ m)
4445
{
4546
init_degrees(m);
4647
init_fabs(m);
48+
init_fix(m);
4749
init_fmax(m);
4850
init_fmin(m);
4951
init_fmod(m);

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

Lines changed: 125 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,125 @@
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+
#include <sycl/sycl.hpp>
27+
28+
#include "dpctl4pybind11.hpp"
29+
30+
#include "fix.hpp"
31+
#include "kernels/elementwise_functions/fix.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 "utils/type_dispatch.hpp"
42+
43+
namespace dpnp::extensions::ufunc
44+
{
45+
namespace py = pybind11;
46+
namespace py_int = dpnp::extensions::py_internal;
47+
48+
namespace impl
49+
{
50+
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
51+
namespace td_ns = dpctl::tensor::type_dispatch;
52+
53+
/**
54+
* @brief A factory to define pairs of supported types for which
55+
* sycl::fix<T> function is available.
56+
*
57+
* @tparam T Type of input vector `a` and of result vector `y`.
58+
*/
59+
template <typename T>
60+
struct OutputType
61+
{
62+
using value_type =
63+
typename std::disjunction<td_ns::TypeMapResultEntry<T, sycl::half>,
64+
td_ns::TypeMapResultEntry<T, float>,
65+
td_ns::TypeMapResultEntry<T, double>,
66+
td_ns::DefaultResultEntry<void>>::result_type;
67+
};
68+
69+
using dpnp::kernels::fix::FixFunctor;
70+
71+
template <typename argT,
72+
typename resT = argT,
73+
unsigned int vec_sz = 4,
74+
unsigned int n_vecs = 2,
75+
bool enable_sg_loadstore = true>
76+
using ContigFunctor = ew_cmn_ns::UnaryContigFunctor<argT,
77+
resT,
78+
FixFunctor<argT, resT>,
79+
vec_sz,
80+
n_vecs,
81+
enable_sg_loadstore>;
82+
83+
template <typename argTy, typename resTy, typename IndexerT>
84+
using StridedFunctor = ew_cmn_ns::
85+
UnaryStridedFunctor<argTy, resTy, IndexerT, FixFunctor<argTy, resTy>>;
86+
87+
using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
88+
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
89+
90+
static unary_contig_impl_fn_ptr_t fix_contig_dispatch_vector[td_ns::num_types];
91+
static int fix_output_typeid_vector[td_ns::num_types];
92+
static unary_strided_impl_fn_ptr_t
93+
fix_strided_dispatch_vector[td_ns::num_types];
94+
95+
MACRO_POPULATE_DISPATCH_VECTORS(fix);
96+
} // namespace impl
97+
98+
void init_fix(py::module_ m)
99+
{
100+
using arrayT = dpctl::tensor::usm_ndarray;
101+
using event_vecT = std::vector<sycl::event>;
102+
{
103+
impl::populate_fix_dispatch_vectors();
104+
using impl::fix_contig_dispatch_vector;
105+
using impl::fix_output_typeid_vector;
106+
using impl::fix_strided_dispatch_vector;
107+
108+
auto fix_pyapi = [&](const arrayT &src, const arrayT &dst,
109+
sycl::queue &exec_q,
110+
const event_vecT &depends = {}) {
111+
return py_int::py_unary_ufunc(
112+
src, dst, exec_q, depends, fix_output_typeid_vector,
113+
fix_contig_dispatch_vector, fix_strided_dispatch_vector);
114+
};
115+
m.def("_fix", fix_pyapi, "", py::arg("src"), py::arg("dst"),
116+
py::arg("sycl_queue"), py::arg("depends") = py::list());
117+
118+
auto fix_result_type_pyapi = [&](const py::dtype &dtype) {
119+
return py_int::py_unary_ufunc_result_type(dtype,
120+
fix_output_typeid_vector);
121+
};
122+
m.def("_fix_result_type", fix_result_type_pyapi);
123+
}
124+
}
125+
} // namespace dpnp::extensions::ufunc

‎dpnp/backend/extensions/ufunc/elementwise_functions/fix.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_fix(py::module_ m);
35+
} // namespace dpnp::extensions::ufunc

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

Lines changed: 49 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,49 @@
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+
namespace dpnp::kernels::fix
31+
{
32+
template <typename argT, typename resT>
33+
struct FixFunctor
34+
{
35+
// is function constant for given argT
36+
using is_constant = typename std::false_type;
37+
// constant value, if constant
38+
// constexpr resT constant_value = resT{};
39+
// is function defined for sycl::vec
40+
using supports_vec = typename std::false_type;
41+
// do both argT and resT support subgroup store/load operation
42+
using supports_sg_loadstore = typename std::true_type;
43+
44+
resT operator()(const argT &x) const
45+
{
46+
return (x >= 0.0) ? sycl::floor(x) : sycl::ceil(x);
47+
}
48+
};
49+
} // namespace dpnp::kernels::fix

‎dpnp/dpnp_iface_mathematical.py

Lines changed: 68 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -98,6 +98,7 @@
9898
"divide",
9999
"ediff1d",
100100
"fabs",
101+
"fix",
101102
"floor",
102103
"floor_divide",
103104
"fmax",
@@ -526,6 +527,7 @@ def around(x, /, decimals=0, out=None):
526527
:obj:`dpnp.round` : Equivalent function; see for details.
527528
:obj:`dpnp.ndarray.round` : Equivalent function.
528529
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
530+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
529531
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
530532
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
531533
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
@@ -571,6 +573,8 @@ def around(x, /, decimals=0, out=None):
571573
--------
572574
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
573575
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
576+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
577+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
574578
575579
Examples
576580
--------
@@ -1364,6 +1368,64 @@ def ediff1d(x1, to_end=None, to_begin=None):
13641368
)
13651369

13661370

1371+
_FIX_DOCSTRING = """
1372+
Round to nearest integer towards zero.
1373+
1374+
Round an array of floats element-wise to nearest integer towards zero.
1375+
The rounded values are returned as floats.
1376+
1377+
For full documentation refer to :obj:`numpy.fix`.
1378+
1379+
Parameters
1380+
----------
1381+
x : {dpnp.ndarray, usm_ndarray}
1382+
The array of numbers for which the absolute values are required.
1383+
out : {None, dpnp.ndarray, usm_ndarray}, optional
1384+
Output array to populate.
1385+
Array must have the correct shape and the expected data type.
1386+
Default: ``None``.
1387+
order : {"C", "F", "A", "K"}, optional
1388+
Memory layout of the newly output array, if parameter `out` is ``None``.
1389+
Default: ``"K"``.
1390+
1391+
Returns
1392+
-------
1393+
out : dpnp.ndarray
1394+
An array with the rounded values and with the same dimensions as the input.
1395+
The returned array will have the default floating point data type for the
1396+
device where `a` is allocated.
1397+
If `out` is ``None`` then a float array is returned with the rounded values.
1398+
Otherwise the result is stored there and the return value `out` is
1399+
a reference to that array.
1400+
1401+
See Also
1402+
--------
1403+
:obj:`dpnp.round` : Round to given number of decimals.
1404+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
1405+
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
1406+
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
1407+
:obj:`dpnp.ceil` : Return the ceiling of the input, element-wise.
1408+
1409+
Examples
1410+
--------
1411+
>>> import dpnp as np
1412+
>>> np.fix(np.array(3.14))
1413+
array(3.)
1414+
>>> np.fix(np.array(3))
1415+
array(3.)
1416+
>>> a = np.array([2.1, 2.9, -2.1, -2.9])
1417+
>>> np.fix(a)
1418+
array([ 2., 2., -2., -2.])
1419+
"""
1420+
1421+
fix = DPNPUnaryFunc(
1422+
"fix",
1423+
ufi._fix_result_type,
1424+
ufi._fix,
1425+
_FIX_DOCSTRING,
1426+
)
1427+
1428+
13671429
_FLOOR_DOCSTRING = """
13681430
Returns the floor for each element `x_i` for input array `x`.
13691431
@@ -1398,6 +1460,8 @@ def ediff1d(x1, to_end=None, to_begin=None):
13981460
--------
13991461
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
14001462
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
1463+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
1464+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
14011465
14021466
Notes
14031467
-----
@@ -2806,6 +2870,7 @@ def prod(
28062870
See Also
28072871
--------
28082872
:obj:`dpnp.round` : Evenly round to the given number of decimals.
2873+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
28092874
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
28102875
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
28112876
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
@@ -2861,6 +2926,7 @@ def prod(
28612926
:obj:`dpnp.ndarray.round` : Equivalent function.
28622927
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
28632928
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
2929+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
28642930
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
28652931
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
28662932
@@ -3294,6 +3360,8 @@ def trapz(y1, x1=None, dx=1.0, axis=-1):
32943360
--------
32953361
:obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity.
32963362
:obj:`dpnp.ceil` : Round a number to the nearest integer toward infinity.
3363+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
3364+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
32973365
32983366
Examples
32993367
--------

0 commit comments

Comments
 (0)