24687578389

Committed 20 Apr 2026 08:03PM UTC coverage: 78.429% (-3.6%) from 82.034%

Build # 24687578389

Build Type

push

github

Committed by

web-flow

Commit Message

Migrate `dpctl.tensor` into `dpnp.tensor` (#2856)

This PR migrates the tensor implementation from `dpctl.tensor` into
`dpnp.tensor` making dpnp the primary owner of the Array API-compliant
tensor layer

Major changes:

- Move compiled C++/SYCL extensions (`_tensor_impl,
_tensor_elementwise_impl, _tensor_reductions_impl, _tensor_sorting_impl,
_tensor_accumulation_impl, tensor linalg`) into `dpnp.tensor`
- Move `usm_ndarray`, `compute-follows-data utilities` and tensor
`tests` from dpctl
- Replace all `dpctl.tensor` references with `dpnp.tensor` in
docstrings, error messages and comments
- Remove redundant dpctl.tensor C-API  interface
- Add `tensor.rst` documentation page describing the module, its
relationship to `dpnp.ndarray` and `dpctl` and linking to the `dpctl
0.21.1 API` reference

This simplifies maintenance, reduces cross-project dependencies and
enables independent development and release cycles

Coverage Stats

1573 of 2908 branches covered (54.09%)

Branch coverage included in aggregate %.

6973 of 9803 new or added lines in 203 files covered. (71.13%)

1 existing line in 1 file now uncovered.

26259 of 32579 relevant lines covered (80.6%)

7622.15 hits per line

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

81.42

/dpnp/backend/extensions/lapack/heevd_batch.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.
// - Neither the name of the copyright holder nor the names of its contributors
//   may be used to endorse or promote products derived from this software
//   without specific prior written permission.
//
// 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 <stdexcept>

#include <pybind11/pybind11.h>
#include <pybind11/stl.h>

#include "common_helpers.hpp"
#include "evd_batch_common.hpp"
#include "heevd_batch.hpp"

// utils extension header
#include "ext/common.hpp"

// dpnp tensor headers
#include "utils/type_utils.hpp"

namespace dpnp::extensions::lapack
{
namespace mkl_lapack = oneapi::mkl::lapack;
namespace type_utils = dpnp::tensor::type_utils;

using ext::common::init_dispatch_table;

template <typename T, typename RealT>
static sycl::event heevd_batch_impl(sycl::queue &exec_q,
                                    const oneapi::mkl::job jobz,
                                    const oneapi::mkl::uplo upper_lower,
                                    const std::int64_t batch_size,
                                    const std::int64_t n,
                                    char *in_a,
                                    char *out_w,
                                    const std::vector<sycl::event> &depends)
{
    type_utils::validate_type_for_device<T>(exec_q);
    type_utils::validate_type_for_device<RealT>(exec_q);

    T *a = reinterpret_cast<T *>(in_a);
    RealT *w = reinterpret_cast<RealT *>(out_w);

    const std::int64_t a_size = n * n;
    const std::int64_t w_size = n;

    const std::int64_t lda = std::max<size_t>(1UL, n);

    // Get the number of independent linear streams
    const std::int64_t n_linear_streams =
        (batch_size > 16) ? 4 : ((batch_size > 4 ? 2 : 1));

    const std::int64_t scratchpad_size =
        mkl_lapack::heevd_scratchpad_size<T>(exec_q, jobz, upper_lower, n, lda);

    T *scratchpad = helper::alloc_scratchpad_batch<T>(scratchpad_size,
                                                      n_linear_streams, exec_q);

    // Computation events to manage dependencies for each linear stream
    std::vector<std::vector<sycl::event>> comp_evs(n_linear_streams, depends);

    std::stringstream error_msg;
    std::int64_t info = 0;

    // Release GIL to avoid serialization of host task
    // submissions to the same queue in OneMKL
    py::gil_scoped_release release;

    for (std::int64_t batch_id = 0; batch_id < batch_size; ++batch_id) {
        T *a_batch = a + batch_id * a_size;
        RealT *w_batch = w + batch_id * w_size;

        std::int64_t stream_id = (batch_id % n_linear_streams);

        T *current_scratch_heevd = scratchpad + stream_id * scratchpad_size;

        // Get the event dependencies for the current stream
        const auto &current_dep = comp_evs[stream_id];

        sycl::event heevd_event;
        try {
            heevd_event = mkl_lapack::heevd(
                exec_q,
                jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors
                      // are computed.
                upper_lower, // 'upper_lower == job::upper' means the upper
                             // triangular part of A, or the lower triangular
                             // otherwise
                n,           // The order of the matrix A (0 <= n)
                a_batch,     // Pointer to the square A (n x n)
                             // If 'jobz == job::vec', then on exit it will
                             // contain the eigenvectors of A
                lda, // The leading dimension of A, must be at least max(1, n)
                w_batch, // Pointer to array of size at least n, it will contain
                         // the eigenvalues of A in ascending order
                current_scratch_heevd, // Pointer to scratchpad memory to be
                                       // used by MKL routine for storing
                                       // intermediate results
                scratchpad_size, current_dep);
        } catch (mkl_lapack::exception const &e) {
            error_msg << "Unexpected MKL exception caught during heevd() "
                         "call:\nreason: "
                      << e.what() << "\ninfo: " << e.info();
            info = e.info();
        } catch (sycl::exception const &e) {
            error_msg
                << "Unexpected SYCL exception caught during heevd() call:\n"
                << e.what();
            info = -1;
        }

        // Update the event dependencies for the current stream
        comp_evs[stream_id] = {heevd_event};
    }

    if (info != 0) // an unexpected error occurs
    {
        if (scratchpad != nullptr) {
            dpnp::tensor::alloc_utils::sycl_free_noexcept(scratchpad, exec_q);
        }
        throw std::runtime_error(error_msg.str());
    }

    sycl::event ht_ev = exec_q.submit([&](sycl::handler &cgh) {
        for (const auto &ev : comp_evs) {
            cgh.depends_on(ev);
        }
        auto ctx = exec_q.get_context();
        cgh.host_task([ctx, scratchpad]() {
            dpnp::tensor::alloc_utils::sycl_free_noexcept(scratchpad, ctx);
        });
    });

    return ht_ev;
}

template <typename fnT, typename T, typename RealT>
struct HeevdBatchContigFactory
{
    fnT get()
    {
        if constexpr (types::HeevdTypePairSupportFactory<T,
                                                         RealT>::is_defined) {
            return heevd_batch_impl<T, RealT>;
        }
        else {
            return nullptr;
        }
    }
};

using evd::evd_batch_impl_fn_ptr_t;

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

    static evd_batch_impl_fn_ptr_t
        heevd_batch_dispatch_table[dpnp_td_ns::num_types]
                                  [dpnp_td_ns::num_types];

    {
        init_dispatch_table<evd_batch_impl_fn_ptr_t, HeevdBatchContigFactory>(
            heevd_batch_dispatch_table);

        auto heevd_batch_pyapi =
            [&](sycl::queue &exec_q, const std::int8_t jobz,
                const std::int8_t upper_lower, const arrayT &eig_vecs,
                const arrayT &eig_vals, const event_vecT &depends = {}) {
                return evd::evd_batch_func(exec_q, jobz, upper_lower, eig_vecs,
                                           eig_vals, depends,
                                           heevd_batch_dispatch_table);
            };
        m.def(
            "_heevd_batch", heevd_batch_pyapi,
            "Call `heevd` from OneMKL LAPACK library in a loop to return "
            "the eigenvalues and eigenvectors of a batch of complex Hermitian "
            "matrices",
            py::arg("sycl_queue"), py::arg("jobz"), py::arg("upper_lower"),
            py::arg("eig_vecs"), py::arg("eig_vals"),
            py::arg("depends") = py::list());
    }
}
} // namespace dpnp::extensions::lapack

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	// - Neither the name of the copyright holder nor the names of its contributors
13	// may be used to endorse or promote products derived from this software
14	// without specific prior written permission.
15	//
16	// THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS"
17	// AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE
18	// IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE
19	// ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT HOLDER OR CONTRIBUTORS BE
20	// LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR
21	// CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF
22	// SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR PROFITS; OR BUSINESS
23	// INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF LIABILITY, WHETHER IN
24	// CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING NEGLIGENCE OR OTHERWISE)
25	// ARISING IN ANY WAY OUT OF THE USE OF THIS SOFTWARE, EVEN IF ADVISED OF
26	// THE POSSIBILITY OF SUCH DAMAGE.
27	//*****************************************************************************
28
29	#include <stdexcept>
30
31	#include <pybind11/pybind11.h>
32	#include <pybind11/stl.h>
33
34	#include "common_helpers.hpp"
35	#include "evd_batch_common.hpp"
36	#include "heevd_batch.hpp"
37
38	// utils extension header
39	#include "ext/common.hpp"
40
41	// dpnp tensor headers
42	#include "utils/type_utils.hpp"
43
44	namespace dpnp::extensions::lapack
45	{
46	namespace mkl_lapack = oneapi::mkl::lapack;
47	namespace type_utils = dpnp::tensor::type_utils;
48
49	using ext::common::init_dispatch_table;
50
51	template <typename T, typename RealT>
52	static sycl::event heevd_batch_impl(sycl::queue &exec_q,
53	const oneapi::mkl::job jobz,
54	const oneapi::mkl::uplo upper_lower,
55	const std::int64_t batch_size,
56	const std::int64_t n,
57	char *in_a,
58	char *out_w,
59	const std::vector<sycl::event> &depends)
60	{	24✔
61	type_utils::validate_type_for_device<T>(exec_q);	24✔
62	type_utils::validate_type_for_device<RealT>(exec_q);	24✔
63
64	T a = reinterpret_cast<T >(in_a);	24✔
65	RealT w = reinterpret_cast<RealT >(out_w);	24✔
66
67	const std::int64_t a_size = n * n;	24✔
68	const std::int64_t w_size = n;	24✔
69
70	const std::int64_t lda = std::max<size_t>(1UL, n);	24✔
71
72	// Get the number of independent linear streams
73	const std::int64_t n_linear_streams =	24✔
74	(batch_size > 16) ? 4 : ((batch_size > 4 ? 2 : 1));	24!
75
76	const std::int64_t scratchpad_size =	24✔
77	mkl_lapack::heevd_scratchpad_size<T>(exec_q, jobz, upper_lower, n, lda);	24✔
78
79	T *scratchpad = helper::alloc_scratchpad_batch<T>(scratchpad_size,	24✔
80	n_linear_streams, exec_q);	24✔
81
82	// Computation events to manage dependencies for each linear stream
83	std::vector<std::vector<sycl::event>> comp_evs(n_linear_streams, depends);	24✔
84
85	std::stringstream error_msg;	24✔
86	std::int64_t info = 0;	24✔
87
88	// Release GIL to avoid serialization of host task
89	// submissions to the same queue in OneMKL
90	py::gil_scoped_release release;	24✔
91
92	for (std::int64_t batch_id = 0; batch_id < batch_size; ++batch_id) {	88✔
93	T a_batch = a + batch_id a_size;	64✔
94	RealT w_batch = w + batch_id w_size;	64✔
95
96	std::int64_t stream_id = (batch_id % n_linear_streams);	64✔
97
98	T current_scratch_heevd = scratchpad + stream_id scratchpad_size;	64✔
99
100	// Get the event dependencies for the current stream
101	const auto &current_dep = comp_evs[stream_id];	64✔
102
103	sycl::event heevd_event;	64✔
104	try {	64✔
105	heevd_event = mkl_lapack::heevd(	64✔
106	exec_q,	64✔
107	jobz, // 'jobz == job::vec' means eigenvalues and eigenvectors	64✔
108	// are computed.
109	upper_lower, // 'upper_lower == job::upper' means the upper	64✔
110	// triangular part of A, or the lower triangular
111	// otherwise
112	n, // The order of the matrix A (0 <= n)	64✔
113	a_batch, // Pointer to the square A (n x n)	64✔
114	// If 'jobz == job::vec', then on exit it will
115	// contain the eigenvectors of A
116	lda, // The leading dimension of A, must be at least max(1, n)	64✔
117	w_batch, // Pointer to array of size at least n, it will contain	64✔
118	// the eigenvalues of A in ascending order
119	current_scratch_heevd, // Pointer to scratchpad memory to be	64✔
120	// used by MKL routine for storing
121	// intermediate results
122	scratchpad_size, current_dep);	64✔
123	} catch (mkl_lapack::exception const &e) {	64✔
124	error_msg << "Unexpected MKL exception caught during heevd() "	×
125	"call:\nreason: "	×
126	<< e.what() << "\ninfo: " << e.info();	×
127	info = e.info();	×
128	} catch (sycl::exception const &e) {	×
129	error_msg	×
130	<< "Unexpected SYCL exception caught during heevd() call:\n"	×
131	<< e.what();	×
132	info = -1;	×
133	}	×
134
135	// Update the event dependencies for the current stream
136	comp_evs[stream_id] = {heevd_event};	64✔
137	}	64✔
138
139	if (info != 0) // an unexpected error occurs	24!
140	{	×
141	if (scratchpad != nullptr) {	×
NEW 142	dpnp::tensor::alloc_utils::sycl_free_noexcept(scratchpad, exec_q);	×
143	}	×
144	throw std::runtime_error(error_msg.str());	×
145	}	×
146
147	sycl::event ht_ev = exec_q.submit([&](sycl::handler &cgh) {	24✔
148	for (const auto &ev : comp_evs) {	24✔
149	cgh.depends_on(ev);	24✔
150	}	24✔
151	auto ctx = exec_q.get_context();	24✔
152	cgh.host_task([ctx, scratchpad]() {	24✔
153	dpnp::tensor::alloc_utils::sycl_free_noexcept(scratchpad, ctx);	24✔
154	});	24✔
155	});	24✔
156
157	return ht_ev;	24✔
158	}	24✔
159
160	template <typename fnT, typename T, typename RealT>
161	struct HeevdBatchContigFactory
162	{
163	fnT get()
164	{	784✔
165	if constexpr (types::HeevdTypePairSupportFactory<T,
166	RealT>::is_defined) {	8✔
167	return heevd_batch_impl<T, RealT>;	8✔
168	}
169	else {	776✔
170	return nullptr;	776✔
171	}	776✔
172	}	784✔
173	};
174
175	using evd::evd_batch_impl_fn_ptr_t;
176
177	void init_heevd_batch(py::module_ m)
178	{	4✔
179	using arrayT = dpnp::tensor::usm_ndarray;	4✔
180	using event_vecT = std::vector<sycl::event>;	4✔
181
182	static evd_batch_impl_fn_ptr_t	4✔
183	heevd_batch_dispatch_table[dpnp_td_ns::num_types]	4✔
184	[dpnp_td_ns::num_types];	4✔
185
186	{	4✔
187	init_dispatch_table<evd_batch_impl_fn_ptr_t, HeevdBatchContigFactory>(	4✔
188	heevd_batch_dispatch_table);	4✔
189
190	auto heevd_batch_pyapi =	4✔
191	[&](sycl::queue &exec_q, const std::int8_t jobz,	4✔
192	const std::int8_t upper_lower, const arrayT &eig_vecs,	4✔
193	const arrayT &eig_vals, const event_vecT &depends = {}) {	24✔
194	return evd::evd_batch_func(exec_q, jobz, upper_lower, eig_vecs,	24✔
195	eig_vals, depends,	24✔
196	heevd_batch_dispatch_table);	24✔
197	};	24✔
198	m.def(	4✔
199	"_heevd_batch", heevd_batch_pyapi,	4✔
200	"Call `heevd` from OneMKL LAPACK library in a loop to return "	4✔
201	"the eigenvalues and eigenvectors of a batch of complex Hermitian "	4✔
202	"matrices",	4✔
203	py::arg("sycl_queue"), py::arg("jobz"), py::arg("upper_lower"),	4✔
204	py::arg("eig_vecs"), py::arg("eig_vals"),	4✔
205	py::arg("depends") = py::list());	4✔
206	}	4✔
207	}	4✔
208	} // namespace dpnp::extensions::lapack

IntelPython / dpnp / 24687578389

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