9488542877

Committed 12 Jun 2024 07:26PM UTC coverage: 40.767% (-0.4%) from 41.215%

Build # 9488542877

Build Type

push

github

Committed by

web-flow

Commit Message

Preparation to reuse future common dpctl f/w in functions from `vm` extension (#1868)

* Preparation to reuse common dpctl f/w for VM functions

* PoC to decouple abs implementation to separate source file

* Reuse typedef for function poiter from dpctl.tensor

* Define populating vectors by a separate macro

* Move implementation of utility functions from headers to source to resolve link issues

* Separated implementation of acos function

* Separated implementation of acosh function

* Use function to simplify strides from dpctl tensor headers

* PoC to decouple add implementation to separate source file

* Separated implementation of asin function

* Separated implementation of asinh function

* Separated implementation of atan, atan2, atanh functions

* Resolve issue with calling MKL function for undefined types

* Separated implementation of cbrt, ceil, conj, cos and cosh functions

* Separated implementation of div, exp, exp2, expm1, floor and hypot functions

* Separated implementation of ln, log1p, log2 and log10 functions

* Separated implementation of mul, pow, rint, sin and sinh functions

* Separated implementation of sqr, sqrt, sub, tan, tanh and trunc functions

* Removed unused header with types matrix

* Remove unused functions

* Use passing by reference in unary and binary funcs

Run Details

3817 of 20909 branches covered (18.26%)

Branch coverage included in aggregate %.

1694 of 2096 new or added lines in 40 files covered. (80.82%)

6 existing lines in 1 file now uncovered.

14048 of 22913 relevant lines covered (61.31%)

32411.82 hits per line

Source File
Press 'n' to go to next uncovered line, 'b' for previous

82.5

/dpnp/backend/extensions/vm/mul.cpp

//*****************************************************************************
// Copyright (c) 2024, Intel Corporation
// All rights reserved.
//
// Redistribution and use in source and binary forms, with or without
// modification, are permitted provided that the following conditions are met:
// - Redistributions of source code must retain the above copyright notice,
//   this list of conditions and the following disclaimer.
// - Redistributions in binary form must reproduce the above copyright notice,
//   this list of conditions and the following disclaimer in the documentation
//   and/or other materials provided with the distribution.
//
// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
// THE POSSIBILITY OF SUCH DAMAGE.
//*****************************************************************************

#include <oneapi/mkl.hpp>
#include <sycl/sycl.hpp>

#include "dpctl4pybind11.hpp"

#include "common.hpp"
#include "mul.hpp"

// include a local copy of elementwise common header from dpctl tensor:
// dpctl/tensor/libtensor/source/elementwise_functions/elementwise_functions.hpp
// TODO: replace by including dpctl header once available
#include "../elementwise_functions/elementwise_functions.hpp"

// dpctl tensor headers
#include "kernels/elementwise_functions/common.hpp"
#include "utils/type_dispatch.hpp"
#include "utils/type_utils.hpp"

namespace dpnp::extensions::vm
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;
namespace tu_ns = dpctl::tensor::type_utils;

namespace impl
{
// OneMKL namespace with VM functions
namespace mkl_vm = oneapi::mkl::vm;

/**
 * @brief A factory to define pairs of supported types for which
 * MKL VM library provides support in oneapi::mkl::vm::mul<T> function.
 *
 * @tparam T Type of input vectors `a` and `b` and of result vector `y`.
 */
template <typename T1, typename T2>
struct OutputType
{
    using value_type = typename std::disjunction<
        td_ns::BinaryTypeMapResultEntry<T1,
                                        std::complex<double>,
                                        T2,
                                        std::complex<double>,
                                        std::complex<double>>,
        td_ns::BinaryTypeMapResultEntry<T1,
                                        std::complex<float>,
                                        T2,
                                        std::complex<float>,
                                        std::complex<float>>,
        td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
        td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
        td_ns::DefaultResultEntry<void>>::result_type;
};

template <typename T1, typename T2>
static sycl::event mul_contig_impl(sycl::queue &exec_q,
                                   std::size_t in_n,
                                   const char *in_a,
                                   ssize_t a_offset,
                                   const char *in_b,
                                   ssize_t b_offset,
                                   char *out_y,
                                   ssize_t out_offset,
                                   const std::vector<sycl::event> &depends)
{
    tu_ns::validate_type_for_device<T1>(exec_q);
    tu_ns::validate_type_for_device<T2>(exec_q);

    if ((a_offset != 0) || (b_offset != 0) || (out_offset != 0)) {
        throw std::runtime_error("Arrays offsets have to be equals to 0");
    }

    std::int64_t n = static_cast<std::int64_t>(in_n);
    const T1 *a = reinterpret_cast<const T1 *>(in_a);
    const T2 *b = reinterpret_cast<const T2 *>(in_b);

    using resTy = typename OutputType<T1, T2>::value_type;
    resTy *y = reinterpret_cast<resTy *>(out_y);

    return mkl_vm::mul(exec_q,
                       n, // number of elements to be calculated
                       a, // pointer `a` containing 1st input vector of size n
                       b, // pointer `b` containing 2nd input vector of size n
                       y, // pointer `y` to the output vector of size n
                       depends);
}

using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
using ew_cmn_ns::binary_strided_impl_fn_ptr_t;

static int output_typeid_vector[td_ns::num_types][td_ns::num_types];
static binary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types]
                                                         [td_ns::num_types];

MACRO_POPULATE_DISPATCH_TABLES(mul);
} // namespace impl

void init_mul(py::module_ m)
{
    using arrayT = dpctl::tensor::usm_ndarray;
    using event_vecT = std::vector<sycl::event>;

    impl::populate_dispatch_tables();
    using impl::contig_dispatch_vector;
    using impl::output_typeid_vector;

    auto mul_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
                         const arrayT &src2, const arrayT &dst,
                         const event_vecT &depends = {}) {
        return py_int::py_binary_ufunc(
            src1, src2, dst, exec_q, depends, output_typeid_vector,
            contig_dispatch_vector,
            // no support of strided implementation in OneMKL
            td_ns::NullPtrTable<impl::binary_strided_impl_fn_ptr_t>{},
            // no support of C-contig row with broadcasting in OneMKL
            td_ns::NullPtrTable<
                impl::
                    binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},
            td_ns::NullPtrTable<
                impl::
                    binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});
    };
    m.def("_mul", mul_pyapi,
          "Call `mul` function from OneMKL VM library to performs element "
          "by element multiplication of vector `src1` by vector `src2` "
          "to resulting vector `dst`",
          py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),
          py::arg("dst"), py::arg("depends") = py::list());

    auto mul_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,
                                      const arrayT &src2, const arrayT &dst) {
        return py_internal::need_to_call_binary_ufunc(exec_q, src1, src2, dst,
                                                      output_typeid_vector,
                                                      contig_dispatch_vector);
    };
    m.def("_mkl_mul_to_call", mul_need_to_call_pyapi,
          "Check input arguments to answer if `mul` function from "
          "OneMKL VM library can be used",
          py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),
          py::arg("dst"));
}
} // namespace dpnp::extensions::vm

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 <oneapi/mkl.hpp>
27	#include <sycl/sycl.hpp>
28
29	#include "dpctl4pybind11.hpp"
30
31	#include "common.hpp"
32	#include "mul.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	#include "utils/type_utils.hpp"
43
44	namespace dpnp::extensions::vm
45	{
46	namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
47	namespace py = pybind11;
48	namespace py_int = dpnp::extensions::py_internal;
49	namespace td_ns = dpctl::tensor::type_dispatch;
50	namespace tu_ns = dpctl::tensor::type_utils;
51
52	namespace impl
53	{
54	// OneMKL namespace with VM functions
55	namespace mkl_vm = oneapi::mkl::vm;
56
57	/**
58	* @brief A factory to define pairs of supported types for which
59	* MKL VM library provides support in oneapi::mkl::vm::mul<T> function.
60	*
61	* @tparam T Type of input vectors `a` and `b` and of result vector `y`.
62	*/
63	template <typename T1, typename T2>
64	struct OutputType
65	{
66	using value_type = typename std::disjunction<
67	td_ns::BinaryTypeMapResultEntry<T1,
68	std::complex<double>,
69	T2,
70	std::complex<double>,
71	std::complex<double>>,
72	td_ns::BinaryTypeMapResultEntry<T1,
73	std::complex<float>,
74	T2,
75	std::complex<float>,
76	std::complex<float>>,
77	td_ns::BinaryTypeMapResultEntry<T1, double, T2, double, double>,
78	td_ns::BinaryTypeMapResultEntry<T1, float, T2, float, float>,
79	td_ns::DefaultResultEntry<void>>::result_type;
80	};
81
82	template <typename T1, typename T2>
83	static sycl::event mul_contig_impl(sycl::queue &exec_q,
84	std::size_t in_n,
85	const char *in_a,
86	ssize_t a_offset,
87	const char *in_b,
88	ssize_t b_offset,
89	char *out_y,
90	ssize_t out_offset,
91	const std::vector<sycl::event> &depends)
92	{	3,128✔
93	tu_ns::validate_type_for_device<T1>(exec_q);	3,128✔
94	tu_ns::validate_type_for_device<T2>(exec_q);	3,128✔
95
96	if ((a_offset != 0) \|\| (b_offset != 0) \|\| (out_offset != 0)) {	3,128!
NEW 97	throw std::runtime_error("Arrays offsets have to be equals to 0");	×
NEW 98	}	×
99
100	std::int64_t n = static_cast<std::int64_t>(in_n);	3,128✔
101	const T1 a = reinterpret_cast<const T1 >(in_a);	3,128✔
102	const T2 b = reinterpret_cast<const T2 >(in_b);	3,128✔
103
104	using resTy = typename OutputType<T1, T2>::value_type;	3,128✔
105	resTy y = reinterpret_cast<resTy >(out_y);	3,128✔
106
107	return mkl_vm::mul(exec_q,	3,128✔
108	n, // number of elements to be calculated	3,128✔
109	a, // pointer `a` containing 1st input vector of size n	3,128✔
110	b, // pointer `b` containing 2nd input vector of size n	3,128✔
111	y, // pointer `y` to the output vector of size n	3,128✔
112	depends);	3,128✔
113	}	3,128✔
114
115	using ew_cmn_ns::binary_contig_impl_fn_ptr_t;
116	using ew_cmn_ns::binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t;
117	using ew_cmn_ns::binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t;
118	using ew_cmn_ns::binary_strided_impl_fn_ptr_t;
119
120	static int output_typeid_vector[td_ns::num_types][td_ns::num_types];
121	static binary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types]
122	[td_ns::num_types];
123
124	MACRO_POPULATE_DISPATCH_TABLES(mul);
125	} // namespace impl
126
127	void init_mul(py::module_ m)
128	{	1✔
129	using arrayT = dpctl::tensor::usm_ndarray;	1✔
130	using event_vecT = std::vector<sycl::event>;	1✔
131
132	impl::populate_dispatch_tables();	1✔
133	using impl::contig_dispatch_vector;	1✔
134	using impl::output_typeid_vector;	1✔
135
136	auto mul_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,	1✔
137	const arrayT &src2, const arrayT &dst,	1✔
138	const event_vecT &depends = {}) {	3,128✔
139	return py_int::py_binary_ufunc(	3,128✔
140	src1, src2, dst, exec_q, depends, output_typeid_vector,	3,128✔
141	contig_dispatch_vector,	3,128✔
142	// no support of strided implementation in OneMKL
143	td_ns::NullPtrTable<impl::binary_strided_impl_fn_ptr_t>{},	3,128✔
144	// no support of C-contig row with broadcasting in OneMKL
145	td_ns::NullPtrTable<	3,128✔
146	impl::	3,128✔
147	binary_contig_matrix_contig_row_broadcast_impl_fn_ptr_t>{},	3,128✔
148	td_ns::NullPtrTable<	3,128✔
149	impl::	3,128✔
150	binary_contig_row_contig_matrix_broadcast_impl_fn_ptr_t>{});	3,128✔
151	};	3,128✔
152	m.def("_mul", mul_pyapi,	1✔
153	"Call `mul` function from OneMKL VM library to performs element "	1✔
154	"by element multiplication of vector `src1` by vector `src2` "	1✔
155	"to resulting vector `dst`",	1✔
156	py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),	1✔
157	py::arg("dst"), py::arg("depends") = py::list());	1✔
158
159	auto mul_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src1,	1✔
160	const arrayT &src2, const arrayT &dst) {	21,690✔
161	return py_internal::need_to_call_binary_ufunc(exec_q, src1, src2, dst,	21,690✔
162	output_typeid_vector,	21,690✔
163	contig_dispatch_vector);	21,690✔
164	};	21,690✔
165	m.def("_mkl_mul_to_call", mul_need_to_call_pyapi,	1✔
166	"Check input arguments to answer if `mul` function from "	1✔
167	"OneMKL VM library can be used",	1✔
168	py::arg("sycl_queue"), py::arg("src1"), py::arg("src2"),	1✔
169	py::arg("dst"));	1✔
170	}	1✔
171	} // namespace dpnp::extensions::vm

IntelPython / dpnp / 9488542877

Source File Press 'n' to go to next uncovered line, 'b' for previous

Source File
Press 'n' to go to next uncovered line, 'b' for previous