11561564489

Committed 28 Oct 2024 07:52PM UTC coverage: 61.533% (-2.7%) from 64.195%

Build # 11561564489

Build Type

Pull #2027

github

Committed by

web-flow

Commit Message

Merge 406dd47ad into 1284f6d00

Pull Request Pull Request #2027: Implementation of histogram with sycl kernel

Run Details

3781 of 9528 branches covered (39.68%)

Branch coverage included in aggregate %.

379 of 703 new or added lines in 8 files covered. (53.91%)

1 existing line in 1 file now uncovered.

14124 of 19570 relevant lines covered (72.17%)

17226.25 hits per line

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

27.54

/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"
// clang-format on

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)
    {
        sycl::atomic_ref<T, Order, Scope> lh(lhs);
        lh += value;
    }
};

template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
struct AtomicOp<std::complex<T>, Order, Scope>
{
    static void add(std::complex<T> &lhs, const std::complex<T> value)
    {
        T *_lhs = reinterpret_cast<T(&)[2]>(lhs);
        const T *_val = reinterpret_cast<const T(&)[2]>(value);
        sycl::atomic_ref<T, Order, Scope> lh0(_lhs[0]);
        lh0 += _val[0];
        sycl::atomic_ref<T, Order, Scope> lh1(_lhs[1]);
        lh1 += _val[1];
    }
};

template <typename T>
struct Less
{
    bool operator()(const T &lhs, const T &rhs) const
    {
        return std::less{}(lhs, rhs);
    }
};

template <typename T>
struct Less<std::complex<T>>
{
    bool operator()(const std::complex<T> &lhs,
                    const std::complex<T> &rhs) const
    {
        return dpctl::tensor::math_utils::less_complex(lhs, rhs);
    }
};

template <typename T>
struct IsNan
{
    static bool isnan(const T &v)
    {
        if constexpr (std::is_floating_point_v<T>) {
            return sycl::isnan(v);
        }

        return false;
    }
};

template <typename T>
struct IsNan<std::complex<T>>
{
    static bool isnan(const std::complex<T> &v)
    {
        T real1 = std::real(v);
        T imag1 = std::imag(v);
        return sycl::isnan(real1) || sycl::isnan(imag1);
    }
};

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 <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	// clang-format on
38
39	namespace statistics
40	{
41	namespace common
42	{
43
44	template <typename N, typename D>
45	constexpr auto CeilDiv(N n, D d)
46	{	401✔
47	return (n + d - 1) / d;	401✔
48	}	401✔
49
50	template <typename N, typename D>
51	constexpr auto Align(N n, D d)
52	{	135✔
53	return CeilDiv(n, d) * d;	135✔
54	}	135✔
55
56	template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
57	struct AtomicOp
58	{
59	static void add(T &lhs, const T value)
NEW 60	{	×
NEW 61	sycl::atomic_ref<T, Order, Scope> lh(lhs);	×
NEW 62	lh += value;	×
NEW 63	}	×
64	};
65
66	template <typename T, sycl::memory_order Order, sycl::memory_scope Scope>
67	struct AtomicOp<std::complex<T>, Order, Scope>
68	{
69	static void add(std::complex<T> &lhs, const std::complex<T> value)
NEW 70	{	×
NEW 71	T *_lhs = reinterpret_cast<T(&)[2]>(lhs);	×
NEW 72	const T *_val = reinterpret_cast<const T(&)[2]>(value);	×
NEW 73	sycl::atomic_ref<T, Order, Scope> lh0(_lhs[0]);	×
NEW 74	lh0 += _val[0];	×
NEW 75	sycl::atomic_ref<T, Order, Scope> lh1(_lhs[1]);	×
NEW 76	lh1 += _val[1];	×
NEW 77	}	×
78	};
79
80	template <typename T>
81	struct Less
82	{
83	bool operator()(const T &lhs, const T &rhs) const
NEW 84	{	×
NEW 85	return std::less{}(lhs, rhs);	×
NEW 86	}	×
87	};
88
89	template <typename T>
90	struct Less<std::complex<T>>
91	{
92	bool operator()(const std::complex<T> &lhs,
93	const std::complex<T> &rhs) const
NEW 94	{	×
NEW 95	return dpctl::tensor::math_utils::less_complex(lhs, rhs);	×
NEW 96	}	×
97	};
98
99	template <typename T>
100	struct IsNan
101	{
102	static bool isnan(const T &v)
NEW 103	{	×
NEW 104	if constexpr (std::is_floating_point_v<T>) {	×
NEW 105	return sycl::isnan(v);	×
NEW 106	}	×
107
NEW 108	return false;	×
NEW 109	}	×
110	};
111
112	template <typename T>
113	struct IsNan<std::complex<T>>
114	{
115	static bool isnan(const std::complex<T> &v)
NEW 116	{	×
NEW 117	T real1 = std::real(v);	×
NEW 118	T imag1 = std::imag(v);	×
NEW 119	return sycl::isnan(real1) \|\| sycl::isnan(imag1);	×
NEW 120	}	×
121	};
122
123	size_t get_max_local_size(const sycl::device &device);
124	size_t get_max_local_size(const sycl::device &device,
125	int cpu_local_size_limit,
126	int gpu_local_size_limit);
127
128	inline size_t get_max_local_size(const sycl::queue &queue)
NEW 129	{	×
NEW 130	return get_max_local_size(queue.get_device());	×
NEW 131	}	×
132
133	inline size_t get_max_local_size(const sycl::queue &queue,
134	int cpu_local_size_limit,
135	int gpu_local_size_limit)
NEW 136	{	×
NEW 137	return get_max_local_size(queue.get_device(), cpu_local_size_limit,	×
NEW 138	gpu_local_size_limit);	×
NEW 139	}	×
140
141	size_t get_local_mem_size_in_bytes(const sycl::device &device);
142	size_t get_local_mem_size_in_bytes(const sycl::device &device, size_t reserve);
143
144	inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue)
NEW 145	{	×
NEW 146	return get_local_mem_size_in_bytes(queue.get_device());	×
NEW 147	}	×
148
149	inline size_t get_local_mem_size_in_bytes(const sycl::queue &queue,
150	size_t reserve)
NEW 151	{	×
NEW 152	return get_local_mem_size_in_bytes(queue.get_device(), reserve);	×
NEW 153	}	×
154
155	template <typename T>
156	size_t get_local_mem_size_in_items(const sycl::device &device)
157	{	135✔
158	return get_local_mem_size_in_bytes(device) / sizeof(T);	135✔
159	}	135✔
160
161	template <typename T>
162	size_t get_local_mem_size_in_items(const sycl::device &device, size_t reserve)
163	{
164	return get_local_mem_size_in_bytes(device, sizeof(T) * reserve) / sizeof(T);
165	}
166
167	template <int Dims>
168	sycl::nd_range<Dims> make_ndrange(const sycl::range<Dims> &global_range,
169	const sycl::range<Dims> &local_range,
170	const sycl::range<Dims> &work_per_item)
171	{	135✔
172	sycl::range<Dims> aligned_global_range;	135✔
173
174	for (int i = 0; i < Dims; ++i) {	270✔
175	aligned_global_range[i] =	135✔
176	Align(CeilDiv(global_range[i], work_per_item[i]), local_range[i]);	135✔
177	}	135✔
178
179	return sycl::nd_range<Dims>(aligned_global_range, local_range);	135✔
180	}	135✔
181
182	sycl::nd_range<1>
183	make_ndrange(size_t global_size, size_t local_range, size_t work_per_item);
184
185	// This function is a copy from dpctl because it is not available in the public
186	// headers of dpctl.
187	pybind11::dtype dtype_from_typenum(int dst_typenum);
188
189	} // namespace common
190	} // namespace statistics

IntelPython / dpnp / 11561564489

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