Skip to content

Commit cab701c

Browse files
committed
use dpctl.tensor.round functions in dpnp
1 parent 07e5f86 commit cab701c

12 files changed

+218
-98
lines changed

dpnp/backend/extensions/vm/round.hpp

+78
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,78 @@
1+
//*****************************************************************************
2+
// Copyright (c) 2023, 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 <CL/sycl.hpp>
29+
30+
#include "common.hpp"
31+
#include "types_matrix.hpp"
32+
33+
namespace dpnp
34+
{
35+
namespace backend
36+
{
37+
namespace ext
38+
{
39+
namespace vm
40+
{
41+
template <typename T>
42+
sycl::event round_contig_impl(sycl::queue exec_q,
43+
const std::int64_t n,
44+
const char *in_a,
45+
char *out_y,
46+
const std::vector<sycl::event> &depends)
47+
{
48+
type_utils::validate_type_for_device<T>(exec_q);
49+
50+
const T *a = reinterpret_cast<const T *>(in_a);
51+
T *y = reinterpret_cast<T *>(out_y);
52+
53+
return mkl_vm::round(exec_q,
54+
n, // number of elements to be calculated
55+
a, // pointer `a` containing input vector of size n
56+
y, // pointer `y` to the output vector of size n
57+
depends);
58+
}
59+
60+
template <typename fnT, typename T>
61+
struct RoundContigFactory
62+
{
63+
fnT get()
64+
{
65+
if constexpr (std::is_same_v<
66+
typename types::RoundOutputType<T>::value_type, void>)
67+
{
68+
return nullptr;
69+
}
70+
else {
71+
return round_contig_impl<T>;
72+
}
73+
}
74+
};
75+
} // namespace vm
76+
} // namespace ext
77+
} // namespace backend
78+
} // namespace dpnp

dpnp/backend/extensions/vm/types_matrix.hpp

+15
Original file line numberDiff line numberDiff line change
@@ -136,6 +136,21 @@ struct LnOutputType
136136
dpctl_td_ns::DefaultResultEntry<void>>::result_type;
137137
};
138138

139+
/**
140+
* @brief A factory to define pairs of supported types for which
141+
* MKL VM library provides support in oneapi::mkl::vm::round<T> function.
142+
*
143+
* @tparam T Type of input vector `a` and of result vector `y`.
144+
*/
145+
template <typename T>
146+
struct RoundOutputType
147+
{
148+
using value_type = typename std::disjunction<
149+
dpctl_td_ns::TypeMapResultEntry<T, double, double>,
150+
dpctl_td_ns::TypeMapResultEntry<T, float, float>,
151+
dpctl_td_ns::DefaultResultEntry<void>>::result_type;
152+
};
153+
139154
/**
140155
* @brief A factory to define pairs of supported types for which
141156
* MKL VM library provides support in oneapi::mkl::vm::sin<T> function.

dpnp/backend/extensions/vm/vm_py.cpp

+30
Original file line numberDiff line numberDiff line change
@@ -36,6 +36,7 @@
3636
#include "div.hpp"
3737
#include "floor.hpp"
3838
#include "ln.hpp"
39+
#include "round.hpp"
3940
#include "sin.hpp"
4041
#include "sqr.hpp"
4142
#include "sqrt.hpp"
@@ -54,6 +55,7 @@ static unary_impl_fn_ptr_t ceil_dispatch_vector[dpctl_td_ns::num_types];
5455
static unary_impl_fn_ptr_t cos_dispatch_vector[dpctl_td_ns::num_types];
5556
static unary_impl_fn_ptr_t floor_dispatch_vector[dpctl_td_ns::num_types];
5657
static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types];
58+
static unary_impl_fn_ptr_t round_dispatch_vector[dpctl_td_ns::num_types];
5759
static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types];
5860
static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types];
5961
static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types];
@@ -206,6 +208,34 @@ PYBIND11_MODULE(_vm_impl, m)
206208
py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
207209
}
208210

211+
// UnaryUfunc: ==== Round(x) ====
212+
{
213+
vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
214+
vm_ext::RoundContigFactory>(
215+
round_dispatch_vector);
216+
217+
auto round_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst,
218+
const event_vecT &depends = {}) {
219+
return vm_ext::unary_ufunc(exec_q, src, dst, depends,
220+
round_dispatch_vector);
221+
};
222+
m.def("_round", round_pyapi,
223+
"Call `round` function from OneMKL VM library to compute "
224+
"the rounded value of vector elements",
225+
py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
226+
py::arg("depends") = py::list());
227+
228+
auto round_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src,
229+
arrayT dst) {
230+
return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst,
231+
round_dispatch_vector);
232+
};
233+
m.def("_mkl_round_to_call", round_need_to_call_pyapi,
234+
"Check input arguments to answer if `round` function from "
235+
"OneMKL VM library can be used",
236+
py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
237+
}
238+
209239
// UnaryUfunc: ==== Sin(x) ====
210240
{
211241
vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,

dpnp/backend/include/dpnp_iface_fptr.hpp

-2
Original file line numberDiff line numberDiff line change
@@ -101,8 +101,6 @@ enum class DPNPFuncName : size_t
101101
DPNP_FN_ARGSORT_EXT, /**< Used in numpy.argsort() impl, requires extra
102102
parameters */
103103
DPNP_FN_AROUND, /**< Used in numpy.around() impl */
104-
DPNP_FN_AROUND_EXT, /**< Used in numpy.around() impl, requires extra
105-
parameters */
106104
DPNP_FN_ASTYPE, /**< Used in numpy.astype() impl */
107105
DPNP_FN_ASTYPE_EXT, /**< Used in numpy.astype() impl, requires extra
108106
parameters */

dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

-18
Original file line numberDiff line numberDiff line change
@@ -109,15 +109,6 @@ template <typename _DataType>
109109
void (*dpnp_around_default_c)(const void *, void *, const size_t, const int) =
110110
dpnp_around_c<_DataType>;
111111

112-
template <typename _DataType>
113-
DPCTLSyclEventRef (*dpnp_around_ext_c)(DPCTLSyclQueueRef,
114-
const void *,
115-
void *,
116-
const size_t,
117-
const int,
118-
const DPCTLEventVectorRef) =
119-
dpnp_around_c<_DataType>;
120-
121112
template <typename _KernelNameSpecialization1,
122113
typename _KernelNameSpecialization2>
123114
class dpnp_elemwise_absolute_c_kernel;
@@ -1201,15 +1192,6 @@ void func_map_init_mathematical(func_map_t &fmap)
12011192
fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = {
12021193
eft_DBL, (void *)dpnp_around_default_c<double>};
12031194

1204-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_INT][eft_INT] = {
1205-
eft_INT, (void *)dpnp_around_ext_c<int32_t>};
1206-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_LNG][eft_LNG] = {
1207-
eft_LNG, (void *)dpnp_around_ext_c<int64_t>};
1208-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_FLT][eft_FLT] = {
1209-
eft_FLT, (void *)dpnp_around_ext_c<float>};
1210-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_DBL][eft_DBL] = {
1211-
eft_DBL, (void *)dpnp_around_ext_c<double>};
1212-
12131195
fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = {
12141196
eft_INT, (void *)dpnp_cross_default_c<int32_t, int32_t, int32_t>};
12151197
fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = {

dpnp/dpnp_algo/dpnp_algo.pxd

-2
Original file line numberDiff line numberDiff line change
@@ -58,8 +58,6 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
5858
DPNP_FN_ARGMIN_EXT
5959
DPNP_FN_ARGSORT
6060
DPNP_FN_ARGSORT_EXT
61-
DPNP_FN_AROUND
62-
DPNP_FN_AROUND_EXT
6361
DPNP_FN_ASTYPE
6462
DPNP_FN_ASTYPE_EXT
6563
DPNP_FN_BITWISE_AND

dpnp/dpnp_algo/dpnp_algo_mathematical.pxi

-36
Original file line numberDiff line numberDiff line change
@@ -38,7 +38,6 @@ and the rest of the library
3838
__all__ += [
3939
"dpnp_absolute",
4040
"dpnp_arctan2",
41-
"dpnp_around",
4241
"dpnp_conjugate",
4342
"dpnp_copysign",
4443
"dpnp_cross",
@@ -76,9 +75,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef,
7675
ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_dpctl.DPCTLSyclQueueRef,
7776
void *, void * , void * , double, size_t, size_t,
7877
const c_dpctl.DPCTLEventVectorRef)
79-
ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_around_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef,
80-
const void * , void * , const size_t, const int,
81-
const c_dpctl.DPCTLEventVectorRef)
8278

8379

8480
cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1):
@@ -124,38 +120,6 @@ cpdef utils.dpnp_descriptor dpnp_arctan2(utils.dpnp_descriptor x1_obj,
124120
return call_fptr_2in_1out_strides(DPNP_FN_ARCTAN2_EXT, x1_obj, x2_obj, dtype, out, where, func_name="arctan2")
125121

126122

127-
cpdef utils.dpnp_descriptor dpnp_around(utils.dpnp_descriptor x1, int decimals):
128-
129-
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(x1.dtype)
130-
131-
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(DPNP_FN_AROUND_EXT, param1_type, param1_type)
132-
133-
x1_obj = x1.get_array()
134-
135-
# ceate result array with type given by FPTR data
136-
cdef shape_type_c result_shape = x1.shape
137-
cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape,
138-
kernel_data.return_type,
139-
None,
140-
device=x1_obj.sycl_device,
141-
usm_type=x1_obj.usm_type,
142-
sycl_queue=x1_obj.sycl_queue)
143-
144-
result_sycl_queue = result.get_array().sycl_queue
145-
146-
cdef c_dpctl.SyclQueue q = <c_dpctl.SyclQueue> result_sycl_queue
147-
cdef c_dpctl.DPCTLSyclQueueRef q_ref = q.get_queue_ref()
148-
149-
cdef ftpr_custom_around_1in_1out_t func = <ftpr_custom_around_1in_1out_t > kernel_data.ptr
150-
151-
cdef c_dpctl.DPCTLSyclEventRef event_ref = func(q_ref, x1.get_data(), result.get_data(), x1.size, decimals, NULL)
152-
153-
with nogil: c_dpctl.DPCTLEvent_WaitAndThrow(event_ref)
154-
c_dpctl.DPCTLEvent_Delete(event_ref)
155-
156-
return result
157-
158-
159123
cpdef utils.dpnp_descriptor dpnp_conjugate(utils.dpnp_descriptor x1):
160124
return call_fptr_1in_1out_strides(DPNP_FN_CONJIGUATE_EXT, x1)
161125

dpnp/dpnp_algo/dpnp_elementwise_common.py

+54
Original file line numberDiff line numberDiff line change
@@ -63,6 +63,7 @@
6363
"dpnp_logical_xor",
6464
"dpnp_multiply",
6565
"dpnp_not_equal",
66+
"dpnp_round",
6667
"dpnp_sin",
6768
"dpnp_sqrt",
6869
"dpnp_square",
@@ -1062,6 +1063,59 @@ def dpnp_not_equal(x1, x2, out=None, order="K"):
10621063
return dpnp_array._create_from_usm_ndarray(res_usm)
10631064

10641065

1066+
_round_docstring = """
1067+
round(x, out=None, order='K')
1068+
1069+
Rounds each element `x_i` of the input array `x` to
1070+
the nearest integer-valued number.
1071+
1072+
Args:
1073+
x (dpnp.ndarray):
1074+
Input array, expected to have numeric data type.
1075+
out ({None, dpnp.ndarray}, optional):
1076+
Output array to populate. Array must have the correct
1077+
shape and the expected data type.
1078+
order ("C","F","A","K", optional): memory layout of the new
1079+
output array, if parameter `out` is `None`.
1080+
Default: "K".
1081+
Return:
1082+
dpnp.ndarray:
1083+
An array containing the element-wise rounded value. The data type
1084+
of the returned array is determined by the Type Promotion Rules.
1085+
"""
1086+
1087+
1088+
def _call_round(src, dst, sycl_queue, depends=None):
1089+
"""A callback to register in UnaryElementwiseFunc class of dpctl.tensor"""
1090+
1091+
if depends is None:
1092+
depends = []
1093+
1094+
if vmi._mkl_round_to_call(sycl_queue, src, dst):
1095+
# call pybind11 extension for round() function from OneMKL VM
1096+
return vmi._round(sycl_queue, src, dst, depends)
1097+
return ti._round(src, dst, sycl_queue, depends)
1098+
1099+
1100+
round_func = UnaryElementwiseFunc(
1101+
"round", ti._round_result_type, _call_round, _round_docstring
1102+
)
1103+
1104+
1105+
def dpnp_round(x, out=None, order="K"):
1106+
"""
1107+
Invokes round() function from pybind11 extension of OneMKL VM if possible.
1108+
1109+
Otherwise fully relies on dpctl.tensor implementation for round() function.
1110+
"""
1111+
# dpctl.tensor only works with usm_ndarray
1112+
x1_usm = dpnp.get_usm_ndarray(x)
1113+
out_usm = None if out is None else dpnp.get_usm_ndarray(out)
1114+
1115+
res_usm = round_func(x1_usm, out=out_usm, order=order)
1116+
return dpnp_array._create_from_usm_ndarray(res_usm)
1117+
1118+
10651119
_sin_docstring = """
10661120
sin(x, out=None, order='K')
10671121
Computes sine for each element `x_i` of input array `x`.

0 commit comments

Comments
 (0)