Skip to content

Commit 5590ceb

Browse files
vlad-perevezentsevvlad-perevezentsev
committed
Merge master into reuse_dpctl_pow
2 parents 816bd93 + a78ae52 commit 5590ceb

28 files changed

+613
-545
lines changed

‎.github/workflows/conda-package.yml

Lines changed: 4 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -43,7 +43,7 @@ jobs:
4343

4444
strategy:
4545
matrix:
46-
python: ['3.8', '3.9', '3.10', '3.11']
46+
python: ['3.9', '3.10', '3.11']
4747
os: [ubuntu-20.04, windows-latest]
4848

4949
runs-on: ${{ matrix.os }}
@@ -120,7 +120,7 @@ jobs:
120120

121121
strategy:
122122
matrix:
123-
python: ['3.8', '3.9', '3.10', '3.11']
123+
python: ['3.9', '3.10', '3.11']
124124
os: [ubuntu-20.04, ubuntu-latest]
125125

126126
continue-on-error: true
@@ -221,7 +221,7 @@ jobs:
221221

222222
strategy:
223223
matrix:
224-
python: ['3.8', '3.9', '3.10', '3.11']
224+
python: ['3.9', '3.10', '3.11']
225225

226226
continue-on-error: true
227227

@@ -353,7 +353,7 @@ jobs:
353353

354354
strategy:
355355
matrix:
356-
python: ['3.8', '3.9', '3.10', '3.11']
356+
python: ['3.9', '3.10', '3.11']
357357
os: [ubuntu-20.04, windows-latest]
358358

359359
runs-on: ${{ matrix.os }}

‎dpnp/backend/extensions/vm/round.hpp

Lines changed: 78 additions & 0 deletions
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::rint(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

Lines changed: 15 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -228,6 +228,21 @@ struct PowOutputType
228228
dpctl_td_ns::DefaultResultEntry<void>>::result_type;
229229
};
230230

231+
/**
232+
* @brief A factory to define pairs of supported types for which
233+
* MKL VM library provides support in oneapi::mkl::vm::rint<T> function.
234+
*
235+
* @tparam T Type of input vector `a` and of result vector `y`.
236+
*/
237+
template <typename T>
238+
struct RoundOutputType
239+
{
240+
using value_type = typename std::disjunction<
241+
dpctl_td_ns::TypeMapResultEntry<T, double, double>,
242+
dpctl_td_ns::TypeMapResultEntry<T, float, float>,
243+
dpctl_td_ns::DefaultResultEntry<void>>::result_type;
244+
};
245+
231246
/**
232247
* @brief A factory to define pairs of supported types for which
233248
* MKL VM library provides support in oneapi::mkl::vm::sin<T> function.

‎dpnp/backend/extensions/vm/vm_py.cpp

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -40,6 +40,7 @@
4040
#include "ln.hpp"
4141
#include "mul.hpp"
4242
#include "pow.hpp"
43+
#include "round.hpp"
4344
#include "sin.hpp"
4445
#include "sqr.hpp"
4546
#include "sqrt.hpp"
@@ -62,6 +63,7 @@ static unary_impl_fn_ptr_t conj_dispatch_vector[dpctl_td_ns::num_types];
6263
static unary_impl_fn_ptr_t ln_dispatch_vector[dpctl_td_ns::num_types];
6364
static binary_impl_fn_ptr_t mul_dispatch_vector[dpctl_td_ns::num_types];
6465
static binary_impl_fn_ptr_t pow_dispatch_vector[dpctl_td_ns::num_types];
66+
static unary_impl_fn_ptr_t round_dispatch_vector[dpctl_td_ns::num_types];
6567
static unary_impl_fn_ptr_t sin_dispatch_vector[dpctl_td_ns::num_types];
6668
static unary_impl_fn_ptr_t sqr_dispatch_vector[dpctl_td_ns::num_types];
6769
static unary_impl_fn_ptr_t sqrt_dispatch_vector[dpctl_td_ns::num_types];
@@ -333,6 +335,34 @@ PYBIND11_MODULE(_vm_impl, m)
333335
py::arg("dst"));
334336
}
335337

338+
// UnaryUfunc: ==== Round(x) ====
339+
{
340+
vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,
341+
vm_ext::RoundContigFactory>(
342+
round_dispatch_vector);
343+
344+
auto round_pyapi = [&](sycl::queue exec_q, arrayT src, arrayT dst,
345+
const event_vecT &depends = {}) {
346+
return vm_ext::unary_ufunc(exec_q, src, dst, depends,
347+
round_dispatch_vector);
348+
};
349+
m.def("_round", round_pyapi,
350+
"Call `rint` function from OneMKL VM library to compute "
351+
"the rounded value of vector elements",
352+
py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
353+
py::arg("depends") = py::list());
354+
355+
auto round_need_to_call_pyapi = [&](sycl::queue exec_q, arrayT src,
356+
arrayT dst) {
357+
return vm_ext::need_to_call_unary_ufunc(exec_q, src, dst,
358+
round_dispatch_vector);
359+
};
360+
m.def("_mkl_round_to_call", round_need_to_call_pyapi,
361+
"Check input arguments to answer if `rint` function from "
362+
"OneMKL VM library can be used",
363+
py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));
364+
}
365+
336366
// UnaryUfunc: ==== Sin(x) ====
337367
{
338368
vm_ext::init_ufunc_dispatch_vector<unary_impl_fn_ptr_t,

‎dpnp/backend/include/dpnp_iface_fptr.hpp

Lines changed: 0 additions & 2 deletions
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

Lines changed: 0 additions & 18 deletions
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;
@@ -1184,15 +1175,6 @@ void func_map_init_mathematical(func_map_t &fmap)
11841175
fmap[DPNPFuncName::DPNP_FN_AROUND][eft_DBL][eft_DBL] = {
11851176
eft_DBL, (void *)dpnp_around_default_c<double>};
11861177

1187-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_INT][eft_INT] = {
1188-
eft_INT, (void *)dpnp_around_ext_c<int32_t>};
1189-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_LNG][eft_LNG] = {
1190-
eft_LNG, (void *)dpnp_around_ext_c<int64_t>};
1191-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_FLT][eft_FLT] = {
1192-
eft_FLT, (void *)dpnp_around_ext_c<float>};
1193-
fmap[DPNPFuncName::DPNP_FN_AROUND_EXT][eft_DBL][eft_DBL] = {
1194-
eft_DBL, (void *)dpnp_around_ext_c<double>};
1195-
11961178
fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_INT] = {
11971179
eft_INT, (void *)dpnp_cross_default_c<int32_t, int32_t, int32_t>};
11981180
fmap[DPNPFuncName::DPNP_FN_CROSS][eft_INT][eft_LNG] = {

‎dpnp/dpnp_algo/dpnp_algo.pxd

Lines changed: 0 additions & 2 deletions
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_CBRT

‎dpnp/dpnp_algo/dpnp_algo.pyx

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -170,7 +170,7 @@ cpdef utils.dpnp_descriptor dpnp_flatten(utils.dpnp_descriptor x1):
170170

171171
cpdef utils.dpnp_descriptor dpnp_init_val(shape, dtype, value):
172172
"""
173-
same as dpnp_full(). TODO remove code dumplication
173+
same as dpnp_full(). TODO remove code duplication
174174
"""
175175
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(dtype)
176176

‎dpnp/dpnp_algo/dpnp_algo_manipulation.pxi

Lines changed: 0 additions & 45 deletions
Original file line numberDiff line numberDiff line change
@@ -38,9 +38,7 @@ and the rest of the library
3838
__all__ += [
3939
"dpnp_atleast_2d",
4040
"dpnp_atleast_3d",
41-
"dpnp_expand_dims",
4241
"dpnp_repeat",
43-
"dpnp_reshape",
4442
]
4543

4644

@@ -104,35 +102,6 @@ cpdef utils.dpnp_descriptor dpnp_atleast_3d(utils.dpnp_descriptor arr):
104102
return arr
105103

106104

107-
cpdef utils.dpnp_descriptor dpnp_expand_dims(utils.dpnp_descriptor in_array, axis):
108-
axis_tuple = utils._object_to_tuple(axis)
109-
result_ndim = len(axis_tuple) + in_array.ndim
110-
111-
if len(axis_tuple) == 0:
112-
axis_ndim = 0
113-
else:
114-
axis_ndim = max(-min(0, min(axis_tuple)), max(0, max(axis_tuple))) + 1
115-
116-
axis_norm = utils._object_to_tuple(utils.normalize_axis(axis_tuple, result_ndim))
117-
118-
if axis_ndim - len(axis_norm) > in_array.ndim:
119-
utils.checker_throw_axis_error("dpnp_expand_dims", "axis", axis, axis_ndim)
120-
121-
if len(axis_norm) > len(set(axis_norm)):
122-
utils.checker_throw_value_error("dpnp_expand_dims", "axis", axis, "no repeated axis")
123-
124-
cdef shape_type_c shape_list
125-
axis_idx = 0
126-
for i in range(result_ndim):
127-
if i in axis_norm:
128-
shape_list.push_back(1)
129-
else:
130-
shape_list.push_back(in_array.shape[axis_idx])
131-
axis_idx = axis_idx + 1
132-
133-
return dpnp_reshape(in_array, shape_list)
134-
135-
136105
cpdef utils.dpnp_descriptor dpnp_repeat(utils.dpnp_descriptor array1, repeats, axes=None):
137106
cdef DPNPFuncType param1_type = dpnp_dtype_to_DPNPFuncType(array1.dtype)
138107

@@ -165,17 +134,3 @@ cpdef utils.dpnp_descriptor dpnp_repeat(utils.dpnp_descriptor array1, repeats, a
165134
c_dpctl.DPCTLEvent_Delete(event_ref)
166135

167136
return result
168-
169-
170-
cpdef utils.dpnp_descriptor dpnp_reshape(utils.dpnp_descriptor array1, newshape, order="C"):
171-
# return dpnp.get_dpnp_descriptor(dpctl.tensor.usm_ndarray(newshape, dtype=numpy.dtype(array1.dtype).name, buffer=array1.get_pyobj()))
172-
# return dpnp.get_dpnp_descriptor(dpctl.tensor.reshape(array1.get_pyobj(), newshape))
173-
array1_obj = array1.get_array()
174-
array_obj = dpctl.tensor.reshape(array1_obj, newshape, order=order)
175-
return dpnp.get_dpnp_descriptor(dpnp_array(array_obj.shape,
176-
buffer=array_obj,
177-
order=order,
178-
device=array1_obj.sycl_device,
179-
usm_type=array1_obj.usm_type,
180-
sycl_queue=array1_obj.sycl_queue),
181-
copy_when_nondefault_queue=False)

‎dpnp/dpnp_algo/dpnp_algo_mathematical.pxi

Lines changed: 0 additions & 36 deletions
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_copysign",
4342
"dpnp_cross",
4443
"dpnp_cumprod",
@@ -71,9 +70,6 @@ ctypedef c_dpctl.DPCTLSyclEventRef(*fptr_1in_2out_t)(c_dpctl.DPCTLSyclQueueRef,
7170
ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_trapz_2in_1out_with_2size_t)(c_dpctl.DPCTLSyclQueueRef,
7271
void *, void * , void * , double, size_t, size_t,
7372
const c_dpctl.DPCTLEventVectorRef)
74-
ctypedef c_dpctl.DPCTLSyclEventRef(*ftpr_custom_around_1in_1out_t)(c_dpctl.DPCTLSyclQueueRef,
75-
const void * , void * , const size_t, const int,
76-
const c_dpctl.DPCTLEventVectorRef)
7773

7874

7975
cpdef utils.dpnp_descriptor dpnp_absolute(utils.dpnp_descriptor x1):
@@ -119,38 +115,6 @@ cpdef utils.dpnp_descriptor dpnp_arctan2(utils.dpnp_descriptor x1_obj,
119115
return call_fptr_2in_1out_strides(DPNP_FN_ARCTAN2_EXT, x1_obj, x2_obj, dtype, out, where, func_name="arctan2")
120116

121117

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

0 commit comments

Comments
 (0)