11818691392

Committed 13 Nov 2024 01:52PM UTC coverage: 68.139% (+0.06%) from 68.084%

Build # 11818691392

Build Type

Pull #2142

github

Committed by

web-flow

Commit Message

Merge 33be327bf into 923eb84a5

Pull Request Pull Request #2142: Implementation of bincounts

Run Details

4351 of 9914 branches covered (43.89%)

Branch coverage included in aggregate %.

187 of 222 new or added lines in 7 files covered. (84.23%)

2 existing lines in 2 files now uncovered.

16310 of 20408 relevant lines covered (79.92%)

20905.93 hits per line

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

35.18

/dpnp/backend/extensions/statistics/histogram.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.
//
// 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 <algorithm>
#include <complex>
#include <memory>
#include <string>
#include <type_traits>
#include <unordered_map>
#include <vector>

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

// dpctl tensor headers
#include "dpctl4pybind11.hpp"
#include "utils/type_dispatch.hpp"

#include "histogram.hpp"
#include "histogram_common.hpp"

namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
using dpctl::tensor::usm_ndarray;

using namespace statistics::histogram;
using namespace statistics::common;

namespace
{

template <typename T, typename DataStorage>
struct HistogramEdges
{
    static constexpr bool const sync_after_init = DataStorage::sync_after_init;
    using boundsT = std::tuple<T, T>;

    HistogramEdges(const T *global_data, size_t size, sycl::handler &cgh)
        : data(global_data, sycl::range<1>(size), cgh)
    {
    }

    template <int _Dims>
    void init(const sycl::nd_item<_Dims> &item) const
    {
        data.init(item);
    }

    boundsT get_bounds() const
    {
        auto min = data.get_ptr()[0];
        auto max = data.get_ptr()[data.size() - 1];
        return {min, max};
    }

    template <int _Dims, typename dT>
    size_t get_bin(const sycl::nd_item<_Dims> &,
                   const dT *val,
                   const boundsT &) const
    {
        uint32_t edges_count = data.size();
        uint32_t bins_count = edges_count - 1;
        const auto *bins = data.get_ptr();

        uint32_t bin =
            std::upper_bound(bins, bins + edges_count, val[0], Less<dT>{}) -
            bins - 1;
        bin = std::min(bin, bins_count - 1);

        return bin;
    }

    template <typename dT>
    bool in_bounds(const dT *val, const boundsT &bounds) const
    {
        return check_in_bounds(val[0], std::get<0>(bounds),
                               std::get<1>(bounds));
    }

private:
    DataStorage data;
};

template <typename T>
using CachedEdges = HistogramEdges<T, CachedData<const T, 1>>;

template <typename T>
using UncachedEdges = HistogramEdges<T, UncachedData<const T, 1>>;

template <typename T, typename BinsT, typename HistType = size_t>
struct HistogramF
{
    static sycl::event impl(sycl::queue &exec_q,
                            const void *vin,
                            const void *vbins_edges,
                            const void *vweights,
                            void *vout,
                            const size_t bins_count,
                            const size_t size,
                            const std::vector<sycl::event> &depends)
    {
        const T *in = static_cast<const T *>(vin);
        const BinsT *bins_edges = static_cast<const BinsT *>(vbins_edges);
        const HistType *weights = static_cast<const HistType *>(vweights);
        HistType *out = static_cast<HistType *>(vout);

        auto device = exec_q.get_device();
        const auto local_size = get_max_local_size(device);

        constexpr uint32_t WorkPI = 128; // empirically found number

        const auto nd_range = make_ndrange(size, local_size, WorkPI);

        return exec_q.submit([&](sycl::handler &cgh) {
            cgh.depends_on(depends);
            constexpr uint32_t dims = 1;

            auto dispatch_edges = [&](uint32_t local_mem, const auto &weights,
                                      auto &hist) {
                if (device.is_gpu() && (local_mem >= bins_count + 1)) {
                    auto edges = CachedEdges(bins_edges, bins_count + 1, cgh);
                    submit_histogram(in, size, dims, WorkPI, hist, edges,
                                     weights, nd_range, cgh);
                }
                else {
                    auto edges = UncachedEdges(bins_edges, bins_count + 1, cgh);
                    submit_histogram(in, size, dims, WorkPI, hist, edges,
                                     weights, nd_range, cgh);
                }
            };

            auto dispatch_bins = [&](const auto &weights) {
                const auto local_mem_size =
                    get_local_mem_size_in_items<T>(device);
                if (local_mem_size >= bins_count) {
                    const auto local_hist_count = get_local_hist_copies_count(
                        local_mem_size, local_size, bins_count);

                    auto hist = HistWithLocalCopies<HistType>(
                        out, bins_count, local_hist_count, cgh);
                    const auto free_local_mem = local_mem_size - hist.size();

                    dispatch_edges(free_local_mem, weights, hist);
                }
                else {
                    auto hist = HistGlobalMemory<HistType>(out);
                    auto edges = UncachedEdges(bins_edges, bins_count + 1, cgh);
                    submit_histogram(in, size, dims, WorkPI, hist, edges,
                                     weights, nd_range, cgh);
                }
            };

            if (weights) {
                auto _weights = Weights(weights);
                dispatch_bins(_weights);
            }
            else {
                auto _weights = NoWeights();
                dispatch_bins(_weights);
            }
        });
    }
};

template <typename SampleType, typename HistType>
using HistogramF_ = HistogramF<SampleType, SampleType, HistType>;

} // namespace

using SupportedTypes = std::tuple<std::tuple<uint64_t, int64_t>,
                                  std::tuple<int64_t, int64_t>,
                                  std::tuple<uint64_t, float>,
                                  std::tuple<int64_t, float>,
                                  std::tuple<uint64_t, double>,
                                  std::tuple<int64_t, double>,
                                  std::tuple<uint64_t, std::complex<float>>,
                                  std::tuple<int64_t, std::complex<float>>,
                                  std::tuple<uint64_t, std::complex<double>>,
                                  std::tuple<int64_t, std::complex<double>>,
                                  std::tuple<float, int64_t>,
                                  std::tuple<double, int64_t>,
                                  std::tuple<float, float>,
                                  std::tuple<double, double>,
                                  std::tuple<float, std::complex<float>>,
                                  std::tuple<double, std::complex<double>>,
                                  std::tuple<std::complex<float>, int64_t>,
                                  std::tuple<std::complex<double>, int64_t>,
                                  std::tuple<std::complex<float>, float>,
                                  std::tuple<std::complex<double>, double>>;

Histogram::Histogram() : dispatch_table("sample", "histogram")
{
    dispatch_table.populate_dispatch_table<SupportedTypes, HistogramF_>();
}

std::tuple<sycl::event, sycl::event>
    Histogram::call(const dpctl::tensor::usm_ndarray &sample,
                    const dpctl::tensor::usm_ndarray &bins,
                    std::optional<const dpctl::tensor::usm_ndarray> &weights,
                    dpctl::tensor::usm_ndarray &histogram,
                    const std::vector<sycl::event> &depends)
{
    validate(sample, bins, weights, histogram);

    if (sample.get_size() == 0) {
        return {sycl::event(), sycl::event()};
    }

    const int sample_typenum = sample.get_typenum();
    const int hist_typenum = histogram.get_typenum();

    auto histogram_func = dispatch_table.get(sample_typenum, hist_typenum);

    auto exec_q = sample.get_queue();

    void *weights_ptr =
        weights.has_value() ? weights.value().get_data() : nullptr;

    auto ev =
        histogram_func(exec_q, sample.get_data(), bins.get_data(), weights_ptr,
                       histogram.get_data(), histogram.get_shape(0),
                       sample.get_shape(0), depends);

    sycl::event args_ev;
    if (weights.has_value()) {
        args_ev = dpctl::utils::keep_args_alive(
            exec_q, {sample, bins, weights.value(), histogram}, {ev});
    }
    else {
        args_ev = dpctl::utils::keep_args_alive(
            exec_q, {sample, bins, histogram}, {ev});
    }

    return {args_ev, ev};
}

std::unique_ptr<Histogram> hist;

void statistics::histogram::populate_histogram(py::module_ m)
{
    using namespace std::placeholders;

    hist.reset(new Histogram());

    auto hist_func =
        [histp = hist.get()](
            const dpctl::tensor::usm_ndarray &sample,
            const dpctl::tensor::usm_ndarray &bins,
            std::optional<const dpctl::tensor::usm_ndarray> &weights,
            dpctl::tensor::usm_ndarray &histogram,
            const std::vector<sycl::event> &depends) {
            return histp->call(sample, bins, weights, histogram, depends);
        };

    m.def("histogram", hist_func, "Compute the histogram of a dataset.",
          py::arg("sample"), py::arg("bins"), py::arg("weights"),
          py::arg("histogram"), py::arg("depends") = py::list());

    auto histogram_dtypes = [histp = hist.get()]() {
        return histp->dispatch_table.get_all_supported_types();
    };

    m.def("histogram_dtypes", histogram_dtypes,
          "Get the supported data types for histogram.");
}

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	#include <algorithm>
27	#include <complex>
28	#include <memory>
29	#include <string>
30	#include <type_traits>
31	#include <unordered_map>
32	#include <vector>
33
34	#include <pybind11/pybind11.h>
35	#include <pybind11/stl.h>
36
37	// dpctl tensor headers
38	#include "dpctl4pybind11.hpp"
39	#include "utils/type_dispatch.hpp"
40
41	#include "histogram.hpp"
42	#include "histogram_common.hpp"
43
44	namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
45	using dpctl::tensor::usm_ndarray;
46
47	using namespace statistics::histogram;
48	using namespace statistics::common;
49
50	namespace
51	{
52
53	template <typename T, typename DataStorage>
54	struct HistogramEdges
55	{
56	static constexpr bool const sync_after_init = DataStorage::sync_after_init;
57	using boundsT = std::tuple<T, T>;
58
59	HistogramEdges(const T *global_data, size_t size, sycl::handler &cgh)
60	: data(global_data, sycl::range<1>(size), cgh)	133✔
61	{	133✔
62	}	133✔
63
64	template <int _Dims>
65	void init(const sycl::nd_item<_Dims> &item) const
66	{	×
67	data.init(item);	×
68	}	×
69
70	boundsT get_bounds() const
71	{	×
72	auto min = data.get_ptr()[0];	×
73	auto max = data.get_ptr()[data.size() - 1];	×
74	return {min, max};	×
75	}	×
76
77	template <int _Dims, typename dT>
78	size_t get_bin(const sycl::nd_item<_Dims> &,
79	const dT *val,
80	const boundsT &) const
81	{	×
82	uint32_t edges_count = data.size();	×
83	uint32_t bins_count = edges_count - 1;	×
84	const auto *bins = data.get_ptr();	×
85
86	uint32_t bin =	×
87	std::upper_bound(bins, bins + edges_count, val[0], Less<dT>{}) -	×
88	bins - 1;	×
89	bin = std::min(bin, bins_count - 1);	×
90
91	return bin;	×
92	}	×
93
94	template <typename dT>
95	bool in_bounds(const dT *val, const boundsT &bounds) const
96	{	×
NEW 97	return check_in_bounds(val[0], std::get<0>(bounds),	×
NEW 98	std::get<1>(bounds));	×
UNCOV 99	}	×
100
101	private:
102	DataStorage data;
103	};
104
105	template <typename T>
106	using CachedEdges = HistogramEdges<T, CachedData<const T, 1>>;
107
108	template <typename T>
109	using UncachedEdges = HistogramEdges<T, UncachedData<const T, 1>>;
110
111	template <typename T, typename BinsT, typename HistType = size_t>
112	struct HistogramF
113	{
114	static sycl::event impl(sycl::queue &exec_q,
115	const void *vin,
116	const void *vbins_edges,
117	const void *vweights,
118	void *vout,
119	const size_t bins_count,
120	const size_t size,
121	const std::vector<sycl::event> &depends)
122	{	133✔
123	const T in = static_cast<const T >(vin);	133✔
124	const BinsT bins_edges = static_cast<const BinsT >(vbins_edges);	133✔
125	const HistType weights = static_cast<const HistType >(vweights);	133✔
126	HistType out = static_cast<HistType >(vout);	133✔
127
128	auto device = exec_q.get_device();	133✔
129	const auto local_size = get_max_local_size(device);	133✔
130
131	constexpr uint32_t WorkPI = 128; // empirically found number	133✔
132
133	const auto nd_range = make_ndrange(size, local_size, WorkPI);	133✔
134
135	return exec_q.submit([&](sycl::handler &cgh) {	133✔
136	cgh.depends_on(depends);	133✔
137	constexpr uint32_t dims = 1;	133✔
138
139	auto dispatch_edges = [&](uint32_t local_mem, const auto &weights,	133✔
140	auto &hist) {	133✔
141	if (device.is_gpu() && (local_mem >= bins_count + 1)) {	131!
142	auto edges = CachedEdges(bins_edges, bins_count + 1, cgh);	×
143	submit_histogram(in, size, dims, WorkPI, hist, edges,	×
144	weights, nd_range, cgh);	×
145	}	×
146	else {	131✔
147	auto edges = UncachedEdges(bins_edges, bins_count + 1, cgh);	131✔
148	submit_histogram(in, size, dims, WorkPI, hist, edges,	131✔
149	weights, nd_range, cgh);	131✔
150	}	131✔
151	};	131✔
152
153	auto dispatch_bins = [&](const auto &weights) {	133✔
154	const auto local_mem_size =	133✔
155	get_local_mem_size_in_items<T>(device);	133✔
156	if (local_mem_size >= bins_count) {	133!
157	const auto local_hist_count = get_local_hist_copies_count(	131✔
158	local_mem_size, local_size, bins_count);	131✔
159
160	auto hist = HistWithLocalCopies<HistType>(	131✔
161	out, bins_count, local_hist_count, cgh);	131✔
162	const auto free_local_mem = local_mem_size - hist.size();	131✔
163
164	dispatch_edges(free_local_mem, weights, hist);	131✔
165	}	131✔
166	else {	2✔
167	auto hist = HistGlobalMemory<HistType>(out);	2✔
168	auto edges = UncachedEdges(bins_edges, bins_count + 1, cgh);	2✔
169	submit_histogram(in, size, dims, WorkPI, hist, edges,	2✔
170	weights, nd_range, cgh);	2✔
171	}	2✔
172	};	133✔
173
174	if (weights) {	133!
175	auto _weights = Weights(weights);	51✔
176	dispatch_bins(_weights);	51✔
177	}	51✔
178	else {	82✔
179	auto _weights = NoWeights();	82✔
180	dispatch_bins(_weights);	82✔
181	}	82✔
182	});	133✔
183	}	133✔
184	};
185
186	template <typename SampleType, typename HistType>
187	using HistogramF_ = HistogramF<SampleType, SampleType, HistType>;
188
189	} // namespace
190
191	using SupportedTypes = std::tuple<std::tuple<uint64_t, int64_t>,
192	std::tuple<int64_t, int64_t>,
193	std::tuple<uint64_t, float>,
194	std::tuple<int64_t, float>,
195	std::tuple<uint64_t, double>,
196	std::tuple<int64_t, double>,
197	std::tuple<uint64_t, std::complex<float>>,
198	std::tuple<int64_t, std::complex<float>>,
199	std::tuple<uint64_t, std::complex<double>>,
200	std::tuple<int64_t, std::complex<double>>,
201	std::tuple<float, int64_t>,
202	std::tuple<double, int64_t>,
203	std::tuple<float, float>,
204	std::tuple<double, double>,
205	std::tuple<float, std::complex<float>>,
206	std::tuple<double, std::complex<double>>,
207	std::tuple<std::complex<float>, int64_t>,
208	std::tuple<std::complex<double>, int64_t>,
209	std::tuple<std::complex<float>, float>,
210	std::tuple<std::complex<double>, double>>;
211
212	Histogram::Histogram() : dispatch_table("sample", "histogram")	1✔
213	{	1✔
214	dispatch_table.populate_dispatch_table<SupportedTypes, HistogramF_>();	1✔
215	}	1✔
216
217	std::tuple<sycl::event, sycl::event>
218	Histogram::call(const dpctl::tensor::usm_ndarray &sample,
219	const dpctl::tensor::usm_ndarray &bins,
220	std::optional<const dpctl::tensor::usm_ndarray> &weights,
221	dpctl::tensor::usm_ndarray &histogram,
222	const std::vector<sycl::event> &depends)
223	{	138✔
224	validate(sample, bins, weights, histogram);	138✔
225
226	if (sample.get_size() == 0) {	138✔
227	return {sycl::event(), sycl::event()};	5✔
228	}	5✔
229
230	const int sample_typenum = sample.get_typenum();	133✔
231	const int hist_typenum = histogram.get_typenum();	133✔
232
233	auto histogram_func = dispatch_table.get(sample_typenum, hist_typenum);	133✔
234
235	auto exec_q = sample.get_queue();	133✔
236
237	void *weights_ptr =	133✔
238	weights.has_value() ? weights.value().get_data() : nullptr;	133✔
239
240	auto ev =	133✔
241	histogram_func(exec_q, sample.get_data(), bins.get_data(), weights_ptr,	133✔
242	histogram.get_data(), histogram.get_shape(0),	133✔
243	sample.get_shape(0), depends);	133✔
244
245	sycl::event args_ev;	133✔
246	if (weights.has_value()) {	133✔
247	args_ev = dpctl::utils::keep_args_alive(	51✔
248	exec_q, {sample, bins, weights.value(), histogram}, {ev});	51✔
249	}	51✔
250	else {	82✔
251	args_ev = dpctl::utils::keep_args_alive(	82✔
252	exec_q, {sample, bins, histogram}, {ev});	82✔
253	}	82✔
254
255	return {args_ev, ev};	133✔
256	}	138✔
257
258	std::unique_ptr<Histogram> hist;
259
260	void statistics::histogram::populate_histogram(py::module_ m)
261	{	1✔
262	using namespace std::placeholders;	1✔
263
264	hist.reset(new Histogram());	1✔
265
266	auto hist_func =	1✔
267	[histp = hist.get()](	1✔
268	const dpctl::tensor::usm_ndarray &sample,	1✔
269	const dpctl::tensor::usm_ndarray &bins,	1✔
270	std::optional<const dpctl::tensor::usm_ndarray> &weights,	1✔
271	dpctl::tensor::usm_ndarray &histogram,	1✔
272	const std::vector<sycl::event> &depends) {	138✔
273	return histp->call(sample, bins, weights, histogram, depends);	138✔
274	};	138✔
275
276	m.def("histogram", hist_func, "Compute the histogram of a dataset.",	1✔
277	py::arg("sample"), py::arg("bins"), py::arg("weights"),	1✔
278	py::arg("histogram"), py::arg("depends") = py::list());	1✔
279
280	auto histogram_dtypes = [histp = hist.get()]() {	138✔
281	return histp->dispatch_table.get_all_supported_types();	138✔
282	};	138✔
283
284	m.def("histogram_dtypes", histogram_dtypes,	1✔
285	"Get the supported data types for histogram.");	1✔
286	}	1✔

IntelPython / dpnp / 11818691392

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