Skip to content

Commit 9558a55

Browse files
antonwolfyantonwolfy
authored
Merge 7be9f7b into 1e5ba88 
2 parents 1e5ba88 + 7be9f7b commit 9558a55

File tree

11 files changed

+349
-4
lines changed

11 files changed

+349
-4
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/float_power.cpp
3132
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmax.cpp
3233
${CMAKE_CURRENT_SOURCE_DIR}/elementwise_functions/fmin.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 "float_power.hpp"
3132
#include "fmax.hpp"
3233
#include "fmin.hpp"
@@ -45,6 +46,7 @@ void init_elementwise_functions(py::module_ m)
4546
{
4647
init_degrees(m);
4748
init_fabs(m);
49+
init_fix(m);
4850
init_float_power(m);
4951
init_fmax(m);
5052
init_fmin(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
"float_power",
102103
"floor",
103104
"floor_divide",
@@ -527,6 +528,7 @@ def around(x, /, decimals=0, out=None):
527528
:obj:`dpnp.round` : Equivalent function; see for details.
528529
:obj:`dpnp.ndarray.round` : Equivalent function.
529530
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
531+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
530532
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
531533
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
532534
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
@@ -572,6 +574,8 @@ def around(x, /, decimals=0, out=None):
572574
--------
573575
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
574576
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
577+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
578+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
575579
576580
Examples
577581
--------
@@ -1365,6 +1369,64 @@ def ediff1d(x1, to_end=None, to_begin=None):
13651369
)
13661370

13671371

1372+
_FIX_DOCSTRING = """
1373+
Round to nearest integer towards zero.
1374+
1375+
Round an array of floats element-wise to nearest integer towards zero.
1376+
The rounded values are returned as floats.
1377+
1378+
For full documentation refer to :obj:`numpy.fix`.
1379+
1380+
Parameters
1381+
----------
1382+
x : {dpnp.ndarray, usm_ndarray}
1383+
The array of numbers for which the absolute values are required.
1384+
out : {None, dpnp.ndarray, usm_ndarray}, optional
1385+
Output array to populate.
1386+
Array must have the correct shape and the expected data type.
1387+
Default: ``None``.
1388+
order : {"C", "F", "A", "K"}, optional
1389+
Memory layout of the newly output array, if parameter `out` is ``None``.
1390+
Default: ``"K"``.
1391+
1392+
Returns
1393+
-------
1394+
out : dpnp.ndarray
1395+
An array with the rounded values and with the same dimensions as the input.
1396+
The returned array will have the default floating point data type for the
1397+
device where `a` is allocated.
1398+
If `out` is ``None`` then a float array is returned with the rounded values.
1399+
Otherwise the result is stored there and the return value `out` is
1400+
a reference to that array.
1401+
1402+
See Also
1403+
--------
1404+
:obj:`dpnp.round` : Round to given number of decimals.
1405+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
1406+
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
1407+
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
1408+
:obj:`dpnp.ceil` : Return the ceiling of the input, element-wise.
1409+
1410+
Examples
1411+
--------
1412+
>>> import dpnp as np
1413+
>>> np.fix(np.array(3.14))
1414+
array(3.)
1415+
>>> np.fix(np.array(3))
1416+
array(3.)
1417+
>>> a = np.array([2.1, 2.9, -2.1, -2.9])
1418+
>>> np.fix(a)
1419+
array([ 2., 2., -2., -2.])
1420+
"""
1421+
1422+
fix = DPNPUnaryFunc(
1423+
"fix",
1424+
ufi._fix_result_type,
1425+
ufi._fix,
1426+
_FIX_DOCSTRING,
1427+
)
1428+
1429+
13681430
_FLOAT_POWER_DOCSTRING = """
13691431
Calculates `x1_i` raised to `x2_i` for each element `x1_i` of the input array
13701432
`x1` with the respective element `x2_i` of the input array `x2`.
@@ -1498,6 +1560,8 @@ def ediff1d(x1, to_end=None, to_begin=None):
14981560
--------
14991561
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
15001562
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
1563+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
1564+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
15011565
15021566
Notes
15031567
-----
@@ -2907,6 +2971,7 @@ def prod(
29072971
See Also
29082972
--------
29092973
:obj:`dpnp.round` : Evenly round to the given number of decimals.
2974+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
29102975
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
29112976
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
29122977
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
@@ -2962,6 +3027,7 @@ def prod(
29623027
:obj:`dpnp.ndarray.round` : Equivalent function.
29633028
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
29643029
:obj:`dpnp.ceil` : Compute the ceiling of the input, element-wise.
3030+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
29653031
:obj:`dpnp.floor` : Return the floor of the input, element-wise.
29663032
:obj:`dpnp.trunc` : Return the truncated value of the input, element-wise.
29673033
@@ -3395,6 +3461,8 @@ def trapz(y1, x1=None, dx=1.0, axis=-1):
33953461
--------
33963462
:obj:`dpnp.floor` : Round a number to the nearest integer toward minus infinity.
33973463
:obj:`dpnp.ceil` : Round a number to the nearest integer toward infinity.
3464+
:obj:`dpnp.rint` : Round elements of the array to the nearest integer.
3465+
:obj:`dpnp.fix` : Round to nearest integer towards zero, element-wise.
33983466
33993467
Examples
34003468
--------

‎tests/skipped_tests.tbl

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -244,8 +244,6 @@ tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_to_nan
244244
tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside
245245
tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside_nan_inf
246246

247-
tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix
248-
249247
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_0_{a_shape=(), b_shape=(), shape=(4, 3, 2)}::test_beta
250248
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_1_{a_shape=(), b_shape=(), shape=(3, 2)}::test_beta
251249
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_2_{a_shape=(), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta

‎tests/skipped_tests_gpu.tbl

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -298,8 +298,6 @@ tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_interp_inf_to_nan
298298
tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside
299299
tests/third_party/cupy/math_tests/test_misc.py::TestMisc::test_heaviside_nan_inf
300300

301-
tests/third_party/cupy/math_tests/test_rounding.py::TestRounding::test_fix
302-
303301
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_0_{a_shape=(), b_shape=(), shape=(4, 3, 2)}::test_beta
304302
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_1_{a_shape=(), b_shape=(), shape=(3, 2)}::test_beta
305303
tests/third_party/cupy/random_tests/test_distributions.py::TestDistributionsBeta_param_2_{a_shape=(), b_shape=(3, 2), shape=(4, 3, 2)}::test_beta

0 commit comments

Comments
 (0)