12246242208

Committed 09 Dec 2024 11:30PM UTC coverage: 65.087% (-1.3%) from 66.389%

Build # 12246242208

Build Type

push

github

Committed by

web-flow

Commit Message

Align with numpy 2.2 (#2226)

The PR proposes to align with changes implemented in `numpy 2.2`:
- Update `dpnp.cov` to properly transpose 2d array with `rowvar=False`.
- Handle boolean arrays in `dpnp.insert` as a mask.
- Mute or workaround tests with fail due to numpy issues in
`numpy.insert` and `numpy.cov`.
- Mute tests with are not passing due to new numpy behavior (new
`matvec` and `vecmat` ufuncs, support of 2d array in `trim_zeros`).

`matvec` and `vecmat` ufuncs and `trim_zeros` functions will be
separately.

Run Details

4560 of 11468 branches covered (39.76%)

Branch coverage included in aggregate %.

4 of 4 new or added lines in 2 files covered. (100.0%)

235 existing lines in 10 files now uncovered.

16946 of 21574 relevant lines covered (78.55%)

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

IntelPython / dpnp / 12246242208

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