11850407869

Committed 15 Nov 2024 04:42AM UTC coverage: 66.445% (-1.7%) from 68.163%

Build # 11850407869

Build Type

Pull #2143

github

Committed by

web-flow

Commit Message

Merge a108da776 into 3b235acee

Pull Request Pull Request #2143: Implementing histogramdd

Run Details

4427 of 10872 branches covered (40.72%)

Branch coverage included in aggregate %.

270 of 384 new or added lines in 6 files covered. (70.31%)

4 existing lines in 2 files now uncovered.

16597 of 20769 relevant lines covered (79.91%)

20555.21 hits per line

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

33.33

/dpnp/backend/extensions/statistics/common.hpp

//*****************************************************************************
// 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.
//*****************************************************************************

#pragma once

#include <complex>
#include <pybind11/numpy.h>
#include <pybind11/pybind11.h>

// clang-format off
// math_utils.hpp doesn't include sycl header but uses sycl types
// so sycl.hpp must be included before math_utils.hpp
#include <sycl/sycl.hpp>
#include "utils/math_utils.hpp"
#include "utils/type_utils.hpp"
// clang-format on

namespace dpctl
{
namespace tensor
{
namespace type_utils
{
// Upstream to dpctl
template <class T>
struct is_complex<const std::complex<T>> : std::true_type
{
};

template <typename T>
constexpr bool is_complex_v = is_complex<T>::value;

} // namespace type_utils
} // namespace tensor
} // namespace dpctl

namespace type_utils = dpctl::tensor::type_utils;

namespace statistics
{
namespace common
{

template <typename N, typename D>
constexpr auto CeilDiv(N n, D d)
{
    return (n + d - 1) / d;
}

template <typename N, typename D>
constexpr auto Align(N n, D d)
{
    return CeilDiv(n, d) * d;
}

template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
struct AtomicOp
{
    static void add(T &lhs, const T &value)
    {
        if constexpr (type_utils::is_complex_v<T>) {
            using vT = typename T::value_type;
            vT *_lhs = reinterpret_cast<vT(&)[2]>(lhs);
            const vT *_val = reinterpret_cast<const vT(&)[2]>(value);

            AtomicOp<vT, Order, Scope>::add(_lhs[0], _val[0]);
            AtomicOp<vT, Order, Scope>::add(_lhs[1], _val[1]);
        }
        else {
            sycl::atomic_ref<T, Order, Scope> lh(lhs);
            lh += value;
        }
    }
};

template <typename T>
struct Less
{
    bool operator()(const T &lhs, const T &rhs) const
    {
        if constexpr (type_utils::is_complex_v<T>) {
            return dpctl::tensor::math_utils::less_complex(lhs, rhs);
        }
        else {
            return std::less{}(lhs, rhs);
        }
    }
};

template <typename T>
struct IsNan
{
    static bool isnan(const T &v)
    {
        if constexpr (type_utils::is_complex_v<T>) {
            const auto real1 = std::real(v);
            const auto imag1 = std::imag(v);

            using vT = typename T::value_type;

            return IsNan<vT>::isnan(real1) || IsNan<vT>::isnan(imag1);
        }
        else {
            if constexpr (std::is_floating_point_v<T> ||
                          std::is_same_v<T, sycl::half>) {
                return sycl::isnan(v);
            }
        }

        return false;
    }
};

size_t get_max_local_size(const sycl::device &device);
size_t get_max_local_size(const sycl::device &device,
                          int cpu_local_size_limit,
                          int gpu_local_size_limit);

inline size_t get_max_local_size(const sycl::queue &queue)
{
    return get_max_local_size(queue.get_device());
}

inline size_t get_max_local_size(const sycl::queue &queue,
                                 int cpu_local_size_limit,
                                 int gpu_local_size_limit)
{
    return get_max_local_size(queue.get_device(), cpu_local_size_limit,
                              gpu_local_size_limit);
}

size_t get_local_mem_size_in_bytes(const sycl::device &device);
size_t get_local_mem_size_in_bytes(const sycl::device &device, size_t reserve);

inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue)
{
    return get_local_mem_size_in_bytes(queue.get_device());
}

inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue,
                                          size_t reserve)
{
    return get_local_mem_size_in_bytes(queue.get_device(), reserve);
}

template <typename T>
size_t get_local_mem_size_in_items(const sycl::device &device)
{
    return get_local_mem_size_in_bytes(device) / sizeof(T);
}

template <typename T>
size_t get_local_mem_size_in_items(const sycl::device &device, size_t reserve)
{
    return get_local_mem_size_in_bytes(device, sizeof(T) * reserve) / sizeof(T);
}

template <typename T>
inline size_t get_local_mem_size_in_items(const sycl::queue &queue)
{
    return get_local_mem_size_in_items<T>(queue.get_device());
}

template <typename T>
inline size_t get_local_mem_size_in_items(const sycl::queue &queue,
                                          size_t reserve)
{
    return get_local_mem_size_in_items<T>(queue.get_device(), reserve);
}

template <int Dims>
sycl::nd_range<Dims> make_ndrange(const sycl::range<Dims> &global_range,
                                  const sycl::range<Dims> &local_range,
                                  const sycl::range<Dims> &work_per_item)
{
    sycl::range<Dims> aligned_global_range;

    for (int i = 0; i < Dims; ++i) {
        aligned_global_range[i] =
            Align(CeilDiv(global_range[i], work_per_item[i]), local_range[i]);
    }

    return sycl::nd_range<Dims>(aligned_global_range, local_range);
}

sycl::nd_range<1>
    make_ndrange(size_t global_size, size_t local_range, size_t work_per_item);

// This function is a copy from dpctl because it is not available in the public
// headers of dpctl.
pybind11::dtype dtype_from_typenum(int dst_typenum);

} // namespace common
} // namespace statistics

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	#pragma once
27
28	#include <complex>
29	#include <pybind11/numpy.h>
30	#include <pybind11/pybind11.h>
31
32	// clang-format off
33	// math_utils.hpp doesn't include sycl header but uses sycl types
34	// so sycl.hpp must be included before math_utils.hpp
35	#include <sycl/sycl.hpp>
36	#include "utils/math_utils.hpp"
37	#include "utils/type_utils.hpp"
38	// clang-format on
39
40	namespace dpctl
41	{
42	namespace tensor
43	{
44	namespace type_utils
45	{
46	// Upstream to dpctl
47	template <class T>
48	struct is_complex<const std::complex<T>> : std::true_type
49	{
50	};
51
52	template <typename T>
53	constexpr bool is_complex_v = is_complex<T>::value;
54
55	} // namespace type_utils
56	} // namespace tensor
57	} // namespace dpctl
58
59	namespace type_utils = dpctl::tensor::type_utils;
60
61	namespace statistics
62	{
63	namespace common
64	{
65
66	template <typename N, typename D>
67	constexpr auto CeilDiv(N n, D d)
68	{	3,109✔
69	return (n + d - 1) / d;	3,109✔
70	}	3,109✔
71
72	template <typename N, typename D>
73	constexpr auto Align(N n, D d)
74	{	833✔
75	return CeilDiv(n, d) * d;	833✔
76	}	833✔
77
78	template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
79	struct AtomicOp
80	{
81	static void add(T &lhs, const T &value)
82	{	×
NEW 83	if constexpr (type_utils::is_complex_v<T>) {	×
NEW 84	using vT = typename T::value_type;	×
NEW 85	vT *_lhs = reinterpret_cast<vT(&)[2]>(lhs);	×
NEW 86	const vT *_val = reinterpret_cast<const vT(&)[2]>(value);	×
87
NEW 88	AtomicOp<vT, Order, Scope>::add(_lhs[0], _val[0]);	×
NEW 89	AtomicOp<vT, Order, Scope>::add(_lhs[1], _val[1]);	×
90	}
NEW 91	else {	×
NEW 92	sycl::atomic_ref<T, Order, Scope> lh(lhs);	×
NEW 93	lh += value;	×
NEW 94	}	×
UNCOV 95	}	×
96	};
97
98	template <typename T>
99	struct Less
100	{
101	bool operator()(const T &lhs, const T &rhs) const
102	{	×
NEW 103	if constexpr (type_utils::is_complex_v<T>) {	×
NEW 104	return dpctl::tensor::math_utils::less_complex(lhs, rhs);	×
105	}
NEW 106	else {	×
NEW 107	return std::less{}(lhs, rhs);	×
NEW 108	}	×
UNCOV 109	}	×
110	};
111
112	template <typename T>
113	struct IsNan
114	{
115	static bool isnan(const T &v)
116	{	×
NEW 117	if constexpr (type_utils::is_complex_v<T>) {	×
NEW 118	const auto real1 = std::real(v);	×
NEW 119	const auto imag1 = std::imag(v);	×
120
NEW 121	using vT = typename T::value_type;	×
122
NEW 123	return IsNan<vT>::isnan(real1) \|\| IsNan<vT>::isnan(imag1);	×
124	}
NEW 125	else {	×
126	if constexpr (std::is_floating_point_v<T> \|\|
NEW 127	std::is_same_v<T, sycl::half>) {	×
NEW 128	return sycl::isnan(v);	×
NEW 129	}	×
UNCOV 130	}	×
131
132	return false;	×
133	}	×
134	};
135
136	size_t get_max_local_size(const sycl::device &device);
137	size_t get_max_local_size(const sycl::device &device,
138	int cpu_local_size_limit,
139	int gpu_local_size_limit);
140
141	inline size_t get_max_local_size(const sycl::queue &queue)
142	{	700✔
143	return get_max_local_size(queue.get_device());	700✔
144	}	700✔
145
146	inline size_t get_max_local_size(const sycl::queue &queue,
147	int cpu_local_size_limit,
148	int gpu_local_size_limit)
149	{	×
150	return get_max_local_size(queue.get_device(), cpu_local_size_limit,	×
151	gpu_local_size_limit);	×
152	}	×
153
154	size_t get_local_mem_size_in_bytes(const sycl::device &device);
155	size_t get_local_mem_size_in_bytes(const sycl::device &device, size_t reserve);
156
157	inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue)
158	{	×
159	return get_local_mem_size_in_bytes(queue.get_device());	×
160	}	×
161
162	inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue,
163	size_t reserve)
164	{	×
165	return get_local_mem_size_in_bytes(queue.get_device(), reserve);	×
166	}	×
167
168	template <typename T>
169	size_t get_local_mem_size_in_items(const sycl::device &device)
170	{	833✔
171	return get_local_mem_size_in_bytes(device) / sizeof(T);	833✔
172	}	833✔
173
174	template <typename T>
175	size_t get_local_mem_size_in_items(const sycl::device &device, size_t reserve)
176	{
177	return get_local_mem_size_in_bytes(device, sizeof(T) * reserve) / sizeof(T);
178	}
179
180	template <typename T>
181	inline size_t get_local_mem_size_in_items(const sycl::queue &queue)
182	{	84✔
183	return get_local_mem_size_in_items<T>(queue.get_device());	84✔
184	}	84✔
185
186	template <typename T>
187	inline size_t get_local_mem_size_in_items(const sycl::queue &queue,
188	size_t reserve)
189	{
190	return get_local_mem_size_in_items<T>(queue.get_device(), reserve);
191	}
192
193	template <int Dims>
194	sycl::nd_range<Dims> make_ndrange(const sycl::range<Dims> &global_range,
195	const sycl::range<Dims> &local_range,
196	const sycl::range<Dims> &work_per_item)
197	{	833✔
198	sycl::range<Dims> aligned_global_range;	833✔
199
200	for (int i = 0; i < Dims; ++i) {	1,666✔
201	aligned_global_range[i] =	833✔
202	Align(CeilDiv(global_range[i], work_per_item[i]), local_range[i]);	833✔
203	}	833✔
204
205	return sycl::nd_range<Dims>(aligned_global_range, local_range);	833✔
206	}	833✔
207
208	sycl::nd_range<1>
209	make_ndrange(size_t global_size, size_t local_range, size_t work_per_item);
210
211	// This function is a copy from dpctl because it is not available in the public
212	// headers of dpctl.
213	pybind11::dtype dtype_from_typenum(int dst_typenum);
214
215	} // namespace common
216	} // namespace statistics

IntelPython / dpnp / 11850407869

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