16126000227

Committed 07 Jul 2025 07:24PM UTC coverage: 22.684% (-49.4%) from 72.051%

Build # 16126000227

Build Type

Pull #2519

github

Committed by

web-flow

Commit Message

Merge bd753a3a3 into 624f14f20

Pull Request Pull Request #2519: tmp changes

Coverage Stats

889 of 9756 branches covered (9.11%)

Branch coverage included in aggregate %.

6317 of 22011 relevant lines covered (28.7%)

35.96 hits per line

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

50.0

/dpnp/backend/extensions/vm/exp.cpp

//*****************************************************************************
// Copyright (c) 2024-2025, 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 "exp.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 py = pybind11;
namespace py_int = dpnp::extensions::py_internal;
namespace td_ns = dpctl::tensor::type_dispatch;

namespace impl
{
namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions
namespace tu_ns = dpctl::tensor::type_utils;

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

template <typename T>
static sycl::event exp_contig_impl(sycl::queue &exec_q,
                                   std::size_t in_n,
                                   const char *in_a,
                                   char *out_y,
                                   const std::vector<sycl::event> &depends)
{
    tu_ns::validate_type_for_device<T>(exec_q);

    std::int64_t n = static_cast<std::int64_t>(in_n);
    const T *a = reinterpret_cast<const T *>(in_a);

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

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

using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
using ew_cmn_ns::unary_strided_impl_fn_ptr_t;

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

MACRO_POPULATE_DISPATCH_VECTORS(exp);
} // namespace impl

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

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

    auto exp_pyapi = [&](sycl::queue &exec_q, const arrayT &src,
                         const arrayT &dst, const event_vecT &depends = {}) {
        return py_int::py_unary_ufunc(
            src, dst, exec_q, depends, output_typeid_vector,
            contig_dispatch_vector,
            // no support of strided implementation in OneMKL
            td_ns::NullPtrVector<impl::unary_strided_impl_fn_ptr_t>{});
    };
    m.def("_exp", exp_pyapi,
          "Call `exp` function from OneMKL VM library to compute "
          "the natural (base-e) exponential of vector elements",
          py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),
          py::arg("depends") = py::list());

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

1	//*****************************************************************************
2	// Copyright (c) 2024-2025, 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 "exp.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 py = pybind11;
47	namespace py_int = dpnp::extensions::py_internal;
48	namespace td_ns = dpctl::tensor::type_dispatch;
49
50	namespace impl
51	{
52	namespace ew_cmn_ns = dpctl::tensor::kernels::elementwise_common;
53	namespace mkl_vm = oneapi::mkl::vm; // OneMKL namespace with VM functions
54	namespace tu_ns = dpctl::tensor::type_utils;
55
56	/**
57	* @brief A factory to define pairs of supported types for which
58	* MKL VM library provides support in oneapi::mkl::vm::exp<T> function.
59	*
60	* @tparam T Type of input vector `a` and of result vector `y`.
61	*/
62	template <typename T>
63	struct OutputType
64	{
65	using value_type = typename std::disjunction<
66	td_ns::TypeMapResultEntry<T, std::complex<double>>,
67	td_ns::TypeMapResultEntry<T, std::complex<float>>,
68	td_ns::TypeMapResultEntry<T, double>,
69	td_ns::TypeMapResultEntry<T, float>,
70	td_ns::DefaultResultEntry<void>>::result_type;
71	};
72
73	template <typename T>
74	static sycl::event exp_contig_impl(sycl::queue &exec_q,
75	std::size_t in_n,
76	const char *in_a,
77	char *out_y,
78	const std::vector<sycl::event> &depends)
79	{	×
80	tu_ns::validate_type_for_device<T>(exec_q);	×
81
82	std::int64_t n = static_cast<std::int64_t>(in_n);	×
83	const T a = reinterpret_cast<const T >(in_a);	×
84
85	using resTy = typename OutputType<T>::value_type;	×
86	resTy y = reinterpret_cast<resTy >(out_y);	×
87
88	return mkl_vm::exp(exec_q,	×
89	n, // number of elements to be calculated	×
90	a, // pointer `a` containing input vector of size n	×
91	y, // pointer `y` to the output vector of size n	×
92	depends);	×
93	}	×
94
95	using ew_cmn_ns::unary_contig_impl_fn_ptr_t;
96	using ew_cmn_ns::unary_strided_impl_fn_ptr_t;
97
98	static int output_typeid_vector[td_ns::num_types];
99	static unary_contig_impl_fn_ptr_t contig_dispatch_vector[td_ns::num_types];
100
101	MACRO_POPULATE_DISPATCH_VECTORS(exp);
102	} // namespace impl
103
104	void init_exp(py::module_ m)
105	{	2✔
106	using arrayT = dpctl::tensor::usm_ndarray;	2✔
107	using event_vecT = std::vector<sycl::event>;	2✔
108
109	impl::populate_dispatch_vectors();	2✔
110	using impl::contig_dispatch_vector;	2✔
111	using impl::output_typeid_vector;	2✔
112
113	auto exp_pyapi = [&](sycl::queue &exec_q, const arrayT &src,	2✔
114	const arrayT &dst, const event_vecT &depends = {}) {	2✔
115	return py_int::py_unary_ufunc(	×
116	src, dst, exec_q, depends, output_typeid_vector,	×
117	contig_dispatch_vector,	×
118	// no support of strided implementation in OneMKL
119	td_ns::NullPtrVector<impl::unary_strided_impl_fn_ptr_t>{});	×
120	};	×
121	m.def("_exp", exp_pyapi,	2✔
122	"Call `exp` function from OneMKL VM library to compute "	2✔
123	"the natural (base-e) exponential of vector elements",	2✔
124	py::arg("sycl_queue"), py::arg("src"), py::arg("dst"),	2✔
125	py::arg("depends") = py::list());	2✔
126
127	auto exp_need_to_call_pyapi = [&](sycl::queue &exec_q, const arrayT &src,	2✔
128	const arrayT &dst) {	2✔
129	return py_internal::need_to_call_unary_ufunc(	×
130	exec_q, src, dst, output_typeid_vector, contig_dispatch_vector);	×
131	};	×
132	m.def("_mkl_exp_to_call", exp_need_to_call_pyapi,	2✔
133	"Check input arguments to answer if `exp` function from "	2✔
134	"OneMKL VM library can be used",	2✔
135	py::arg("sycl_queue"), py::arg("src"), py::arg("dst"));	2✔
136	}	2✔
137	} // namespace dpnp::extensions::vm

IntelPython / dpnp / 16126000227

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