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

22.92

/dpnp/backend/extensions/window/kaiser.cpp

//*****************************************************************************
// Copyright (c) 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 "kaiser.hpp"
#include "common.hpp"

#include "utils/output_validation.hpp"
#include "utils/type_dispatch.hpp"
#include "utils/type_utils.hpp"

#include <sycl/sycl.hpp>

#include "../kernels/elementwise_functions/i0.hpp"

namespace dpnp::extensions::window
{
namespace dpctl_td_ns = dpctl::tensor::type_dispatch;

typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &,
                                       char *,
                                       const std::size_t,
                                       const py::object &,
                                       const std::vector<sycl::event> &);

static kaiser_fn_ptr_t kaiser_dispatch_vector[dpctl_td_ns::num_types];

template <typename T>
class KaiserFunctor
{
private:
    T *res = nullptr;
    const std::size_t N;
    const T beta;

public:
    KaiserFunctor(T *res, const std::size_t N, const T beta)
        : res(res), N(N), beta(beta)
    {
    }

    void operator()(sycl::id<1> id) const
    {
        using dpnp::kernels::i0::cyl_bessel_i0;

        const auto i = id.get(0);
        const T alpha = (N - 1) / T(2);
        const T tmp = (i - alpha) / alpha;
        res[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) /
                 cyl_bessel_i0(beta);
    }
};

template <typename T>
sycl::event kaiser_impl(sycl::queue &exec_q,
                        char *result,
                        const std::size_t nelems,
                        const py::object &py_beta,
                        const std::vector<sycl::event> &depends)
{
    dpctl::tensor::type_utils::validate_type_for_device<T>(exec_q);

    T *res = reinterpret_cast<T *>(result);
    const T beta = py::cast<const T>(py_beta);

    sycl::event kaiser_ev = exec_q.submit([&](sycl::handler &cgh) {
        cgh.depends_on(depends);

        using KaiserKernel = KaiserFunctor<T>;
        cgh.parallel_for<KaiserKernel>(sycl::range<1>(nelems),
                                       KaiserKernel(res, nelems, beta));
    });

    return kaiser_ev;
}

template <typename fnT, typename T>
struct KaiserFactory
{
    fnT get()
    {
        if constexpr (std::is_floating_point_v<T>) {
            return kaiser_impl<T>;
        }
        else {
            return nullptr;
        }
    }
};

std::pair<sycl::event, sycl::event>
    py_kaiser(sycl::queue &exec_q,
              const py::object &py_beta,
              const dpctl::tensor::usm_ndarray &result,
              const std::vector<sycl::event> &depends)
{
    auto [nelems, result_typeless_ptr, kaiser_fn] =
        window_fn<kaiser_fn_ptr_t>(exec_q, result, kaiser_dispatch_vector);

    if (nelems == 0) {
        return std::make_pair(sycl::event{}, sycl::event{});
    }

    sycl::event kaiser_ev =
        kaiser_fn(exec_q, result_typeless_ptr, nelems, py_beta, depends);
    sycl::event args_ev =
        dpctl::utils::keep_args_alive(exec_q, {result}, {kaiser_ev});

    return std::make_pair(args_ev, kaiser_ev);
}

void init_kaiser_dispatch_vectors()
{
    init_window_dispatch_vectors<kaiser_fn_ptr_t, KaiserFactory>(
        kaiser_dispatch_vector);
}

} // namespace dpnp::extensions::window

1	//*****************************************************************************
2	// Copyright (c) 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 "kaiser.hpp"
27	#include "common.hpp"
28
29	#include "utils/output_validation.hpp"
30	#include "utils/type_dispatch.hpp"
31	#include "utils/type_utils.hpp"
32
33	#include <sycl/sycl.hpp>
34
35	#include "../kernels/elementwise_functions/i0.hpp"
36
37	namespace dpnp::extensions::window
38	{
39	namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
40
41	typedef sycl::event (*kaiser_fn_ptr_t)(sycl::queue &,
42	char *,
43	const std::size_t,
44	const py::object &,
45	const std::vector<sycl::event> &);
46
47	static kaiser_fn_ptr_t kaiser_dispatch_vector[dpctl_td_ns::num_types];
48
49	template <typename T>
50	class KaiserFunctor
51	{
52	private:
53	T *res = nullptr;
54	const std::size_t N;
55	const T beta;
56
57	public:
58	KaiserFunctor(T *res, const std::size_t N, const T beta)
59	: res(res), N(N), beta(beta)	×
60	{	×
61	}	×
62
63	void operator()(sycl::id<1> id) const
64	{	×
65	using dpnp::kernels::i0::cyl_bessel_i0;	×
66
67	const auto i = id.get(0);	×
68	const T alpha = (N - 1) / T(2);	×
69	const T tmp = (i - alpha) / alpha;	×
70	res[i] = cyl_bessel_i0(beta * sycl::sqrt(1 - tmp * tmp)) /	×
71	cyl_bessel_i0(beta);	×
72	}	×
73	};
74
75	template <typename T>
76	sycl::event kaiser_impl(sycl::queue &exec_q,
77	char *result,
78	const std::size_t nelems,
79	const py::object &py_beta,
80	const std::vector<sycl::event> &depends)
81	{	×
82	dpctl::tensor::type_utils::validate_type_for_device<T>(exec_q);	×
83
84	T res = reinterpret_cast<T >(result);	×
85	const T beta = py::cast<const T>(py_beta);	×
86
87	sycl::event kaiser_ev = exec_q.submit([&](sycl::handler &cgh) {	×
88	cgh.depends_on(depends);	×
89
90	using KaiserKernel = KaiserFunctor<T>;	×
91	cgh.parallel_for<KaiserKernel>(sycl::range<1>(nelems),	×
92	KaiserKernel(res, nelems, beta));	×
93	});	×
94
95	return kaiser_ev;	×
96	}	×
97
98	template <typename fnT, typename T>
99	struct KaiserFactory
100	{
101	fnT get()
102	{	28✔
103	if constexpr (std::is_floating_point_v<T>) {	28✔
104	return kaiser_impl<T>;	4✔
105	}
106	else {	24✔
107	return nullptr;	24✔
108	}	24✔
109	}	28✔
110	};
111
112	std::pair<sycl::event, sycl::event>
113	py_kaiser(sycl::queue &exec_q,
114	const py::object &py_beta,
115	const dpctl::tensor::usm_ndarray &result,
116	const std::vector<sycl::event> &depends)
117	{	×
118	auto [nelems, result_typeless_ptr, kaiser_fn] =	×
119	window_fn<kaiser_fn_ptr_t>(exec_q, result, kaiser_dispatch_vector);	×
120
121	if (nelems == 0) {	×
122	return std::make_pair(sycl::event{}, sycl::event{});	×
123	}	×
124
125	sycl::event kaiser_ev =	×
126	kaiser_fn(exec_q, result_typeless_ptr, nelems, py_beta, depends);	×
127	sycl::event args_ev =	×
128	dpctl::utils::keep_args_alive(exec_q, {result}, {kaiser_ev});	×
129
130	return std::make_pair(args_ev, kaiser_ev);	×
131	}	×
132
133	void init_kaiser_dispatch_vectors()
134	{	2✔
135	init_window_dispatch_vectors<kaiser_fn_ptr_t, KaiserFactory>(	2✔
136	kaiser_dispatch_vector);	2✔
137	}	2✔
138
139	} // namespace dpnp::extensions::window

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