11707166170

Committed 06 Nov 2024 03:53PM UTC coverage: 60.49% (-1.7%) from 62.216%

Build # 11707166170

Build Type

Pull #2143

github

Committed by

web-flow

Commit Message

Merge 061961ff3 into 687b8eacc

Pull Request Pull Request #2143: Implementing histogramdd

Run Details

4052 of 10826 branches covered (37.43%)

Branch coverage included in aggregate %.

264 of 383 new or added lines in 6 files covered. (68.93%)

27 existing lines in 7 files now uncovered.

14882 of 20475 relevant lines covered (72.68%)

16479.47 hits per line

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

54.26

/dpnp/backend/extensions/statistics/histogram_common.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 <limits>
#include <string>
#include <unordered_map>
#include <vector>

#include "dpctl4pybind11.hpp"
#include "utils/memory_overlap.hpp"
#include "utils/output_validation.hpp"
#include "utils/type_dispatch.hpp"

#include <pybind11/pybind11.h>

#include "histogram_common.hpp"

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

namespace statistics
{
using common::CeilDiv;

namespace histogram
{

void validate(const usm_ndarray &sample,
              const usm_ndarray &bins,
              std::optional<const dpctl::tensor::usm_ndarray> &weights,
              const usm_ndarray &histogram)
{
    auto exec_q = sample.get_queue();
    using array_ptr = const usm_ndarray *;

    std::vector<array_ptr> arrays{&sample, &bins, &histogram};
    std::unordered_map<array_ptr, std::string> names = {
        {arrays[0], "sample"}, {arrays[1], "bins"}, {arrays[2], "histogram"}};

    array_ptr weights_ptr = nullptr;

    if (weights.has_value()) {
        weights_ptr = &weights.value();
        arrays.push_back(weights_ptr);
        names.insert({weights_ptr, "weights"});
    }

    auto get_name = [&](const array_ptr &arr) {
        auto name_it = names.find(arr);
        assert(name_it != names.end());

        return "'" + name_it->second + "'";
    };

    dpctl::tensor::validation::CheckWritable::throw_if_not_writable(histogram);

    auto unequal_queue =
        std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {
            return arr->get_queue() != exec_q;
        });

    if (unequal_queue != arrays.cend()) {
        throw py::value_error(
            get_name(*unequal_queue) +
            " parameter has incompatible queue with parameter " +
            get_name(&sample));
    }

    auto non_contig_array =
        std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {
            return !arr->is_c_contiguous();
        });

    if (non_contig_array != arrays.cend()) {
        throw py::value_error(get_name(*non_contig_array) +
                              " parameter is not c-contiguos");
    }

    auto check_overlaping = [&](const array_ptr &first,
                                const array_ptr &second) {
        if (first == nullptr || second == nullptr) {
            return;
        }

        const auto &overlap = dpctl::tensor::overlap::MemoryOverlap();

        if (overlap(*first, *second)) {
            throw py::value_error(get_name(first) +
                                  " has overlapping memory segments with " +
                                  get_name(second));
        }
    };

    check_overlaping(&sample, &histogram);
    check_overlaping(&bins, &histogram);
    check_overlaping(weights_ptr, &histogram);

    if (bins.get_size() < 2) {
        throw py::value_error(get_name(&bins) +
                              " parameter must have at least 2 elements");
    }

    if (histogram.get_size() < 1) {
        throw py::value_error(get_name(&histogram) +
                              " parameter must have at least 1 element");
    }

    if (weights_ptr) {
        if (weights_ptr->get_ndim() != 1) {
            throw py::value_error(
                get_name(weights_ptr) + " parameter must be 1d. Actual " +
                std::to_string(weights_ptr->get_ndim()) + "d");
        }

        auto sample_size = sample.get_shape(0);
        auto weights_size = weights_ptr->get_size();
        if (sample_size != weights_ptr->get_size()) {
            throw py::value_error(
                get_name(&sample) + " size (" + std::to_string(sample_size) +
                ") and " + get_name(weights_ptr) + " size (" +
                std::to_string(weights_size) + ")" + " must match");
        }
    }

    if (sample.get_ndim() > 2) {
        throw py::value_error(
            get_name(&sample) +
            " parameter must have no more than 2 dimensions. Actual " +
            std::to_string(sample.get_ndim()) + "d");
    }

    if (sample.get_ndim() == 1) {
        if (histogram.get_ndim() != 1) {
            throw py::value_error(get_name(&sample) + " parameter is 1d, but " +
                                  get_name(&histogram) + " is " +
                                  std::to_string(histogram.get_ndim()) + "d");
        }

        if (histogram.get_size() != bins.get_size() - 1) {
            auto hist_size = histogram.get_size();
            auto bins_size = bins.get_size();
            throw py::value_error(
                get_name(&histogram) + " parameter and " + get_name(&bins) +
                " parameters shape mismatch. " + get_name(&histogram) +
                " size is " + std::to_string(hist_size) + get_name(&bins) +
                " must have size " + std::to_string(hist_size + 1) +
                " but have " + std::to_string(bins_size));
        }
    }
    else if (sample.get_ndim() == 2) {
        auto sample_count = sample.get_shape(0);
        auto expected_dims = sample.get_shape(1);

        if (histogram.get_ndim() != expected_dims) {
            throw py::value_error(get_name(&sample) + " parameter has shape (" +
                                  std::to_string(sample_count) + ", " +
                                  std::to_string(expected_dims) + ")" +
                                  ", so " + get_name(&histogram) +
                                  " parameter expected to be " +
                                  std::to_string(expected_dims) +
                                  "d. "
                                  "Actual " +
                                  std::to_string(histogram.get_ndim()) + "d");
        }
    }

    int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;
    if (histogram.get_size() > max_hist_size) {
        throw py::value_error(get_name(&histogram) +
                              " parameter size expected to be less than " +
                              std::to_string(max_hist_size) + ". Actual " +
                              std::to_string(histogram.get_size()));
    }

    auto array_types = dpctl_td_ns::usm_ndarray_types();
    auto hist_type = static_cast<typenum_t>(
        array_types.typenum_to_lookup_id(histogram.get_typenum()));
    if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {
        auto device = exec_q.get_device();
        bool _64bit_atomics = device.has(sycl::aspect::atomic64);

        if (!_64bit_atomics) {
            auto device_name = device.get_info<sycl::info::device::name>();
            throw py::value_error(
                get_name(&histogram) +
                " parameter has 64-bit type, but 64-bit atomics " +
                " are not supported for " + device_name);
        }
    }
}

uint32_t get_local_hist_copies_count(uint32_t loc_mem_size_in_items,
                                     uint32_t local_size,
                                     uint32_t hist_size_in_items)
{
    constexpr uint32_t local_copies_limit = 16;
    constexpr uint32_t atomics_per_work_item = 4;

    const uint32_t preferred_local_copies =
        CeilDiv(atomics_per_work_item * local_size, hist_size_in_items);
    const uint32_t local_copies_fit_memory =
        loc_mem_size_in_items / hist_size_in_items;

    uint32_t local_hist_count =
        std::min(preferred_local_copies, local_copies_limit);
    local_hist_count = std::min(local_hist_count, local_copies_fit_memory);

    return local_hist_count;
}

} // namespace histogram
} // 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	#include <algorithm>
27	#include <limits>
28	#include <string>
29	#include <unordered_map>
30	#include <vector>
31
32	#include "dpctl4pybind11.hpp"
33	#include "utils/memory_overlap.hpp"
34	#include "utils/output_validation.hpp"
35	#include "utils/type_dispatch.hpp"
36
37	#include <pybind11/pybind11.h>
38
39	#include "histogram_common.hpp"
40
41	namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
42	using dpctl::tensor::usm_ndarray;
43	using dpctl_td_ns::typenum_t;
44
45	namespace statistics
46	{
47	using common::CeilDiv;
48
49	namespace histogram
50	{
51
52	void validate(const usm_ndarray &sample,
53	const usm_ndarray &bins,
54	std::optional<const dpctl::tensor::usm_ndarray> &weights,
55	const usm_ndarray &histogram)
56	{	756✔
57	auto exec_q = sample.get_queue();	756✔
58	using array_ptr = const usm_ndarray *;	756✔
59
60	std::vector<array_ptr> arrays{&sample, &bins, &histogram};	756✔
61	std::unordered_map<array_ptr, std::string> names = {	756✔
62	{arrays[0], "sample"}, {arrays[1], "bins"}, {arrays[2], "histogram"}};	756✔
63
64	array_ptr weights_ptr = nullptr;	756✔
65
66	if (weights.has_value()) {	756✔
67	weights_ptr = &weights.value();	445✔
68	arrays.push_back(weights_ptr);	445✔
69	names.insert({weights_ptr, "weights"});	445✔
70	}	445✔
71
72	auto get_name = [&](const array_ptr &arr) {	756✔
73	auto name_it = names.find(arr);	×
74	assert(name_it != names.end());	×
75
76	return "'" + name_it->second + "'";	×
77	};	×
78
79	dpctl::tensor::validation::CheckWritable::throw_if_not_writable(histogram);	756✔
80
81	auto unequal_queue =	756✔
82	std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {	2,713✔
83	return arr->get_queue() != exec_q;	2,713✔
84	});	2,713✔
85
86	if (unequal_queue != arrays.cend()) {	756!
87	throw py::value_error(	×
88	get_name(*unequal_queue) +	×
89	" parameter has incompatible queue with parameter " +	×
90	get_name(&sample));	×
91	}	×
92
93	auto non_contig_array =	756✔
94	std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {	2,713✔
95	return !arr->is_c_contiguous();	2,713✔
96	});	2,713✔
97
98	if (non_contig_array != arrays.cend()) {	756!
99	throw py::value_error(get_name(*non_contig_array) +	×
100	" parameter is not c-contiguos");	×
101	}	×
102
103	auto check_overlaping = [&](const array_ptr &first,	756✔
104	const array_ptr &second) {	2,268✔
105	if (first == nullptr \|\| second == nullptr) {	2,268!
106	return;	311✔
107	}	311✔
108
109	const auto &overlap = dpctl::tensor::overlap::MemoryOverlap();	1,957✔
110
111	if (overlap(first, second)) {	1,957!
112	throw py::value_error(get_name(first) +	×
113	" has overlapping memory segments with " +	×
114	get_name(second));	×
115	}	×
116	};	1,957✔
117
118	check_overlaping(&sample, &histogram);	756✔
119	check_overlaping(&bins, &histogram);	756✔
120	check_overlaping(weights_ptr, &histogram);	756✔
121
122	if (bins.get_size() < 2) {	756!
123	throw py::value_error(get_name(&bins) +	×
124	" parameter must have at least 2 elements");	×
125	}	×
126
127	if (histogram.get_size() < 1) {	756!
128	throw py::value_error(get_name(&histogram) +	×
129	" parameter must have at least 1 element");	×
130	}	×
131
132	if (weights_ptr) {	756✔
133	if (weights_ptr->get_ndim() != 1) {	445!
134	throw py::value_error(	×
135	get_name(weights_ptr) + " parameter must be 1d. Actual " +	×
136	std::to_string(weights_ptr->get_ndim()) + "d");	×
137	}	×
138
139	auto sample_size = sample.get_shape(0);	445✔
140	auto weights_size = weights_ptr->get_size();	445✔
141	if (sample_size != weights_ptr->get_size()) {	445!
142	throw py::value_error(	×
143	get_name(&sample) + " size (" + std::to_string(sample_size) +	×
144	") and " + get_name(weights_ptr) + " size (" +	×
145	std::to_string(weights_size) + ")" + " must match");	×
146	}	×
147	}	445✔
148
149	if (sample.get_ndim() > 2) {	756!
150	throw py::value_error(	×
151	get_name(&sample) +	×
152	" parameter must have no more than 2 dimensions. Actual " +	×
153	std::to_string(sample.get_ndim()) + "d");	×
154	}	×
155
156	if (sample.get_ndim() == 1) {	756✔
157	if (histogram.get_ndim() != 1) {	140!
158	throw py::value_error(get_name(&sample) + " parameter is 1d, but " +	×
NEW 159	get_name(&histogram) + " is " +	×
NEW 160	std::to_string(histogram.get_ndim()) + "d");	×
NEW 161	}	×
162
163	if (histogram.get_size() != bins.get_size() - 1) {	140!
NEW 164	auto hist_size = histogram.get_size();	×
NEW 165	auto bins_size = bins.get_size();	×
NEW 166	throw py::value_error(	×
NEW 167	get_name(&histogram) + " parameter and " + get_name(&bins) +	×
NEW 168	" parameters shape mismatch. " + get_name(&histogram) +	×
NEW 169	" size is " + std::to_string(hist_size) + get_name(&bins) +	×
NEW 170	" must have size " + std::to_string(hist_size + 1) +	×
NEW 171	" but have " + std::to_string(bins_size));	×
UNCOV 172	}	×
173	}	140✔
174	else if (sample.get_ndim() == 2) {	616!
175	auto sample_count = sample.get_shape(0);	616✔
176	auto expected_dims = sample.get_shape(1);	616✔
177
178	if (histogram.get_ndim() != expected_dims) {	616!
NEW 179	throw py::value_error(get_name(&sample) + " parameter has shape (" +	×
NEW 180	std::to_string(sample_count) + ", " +	×
NEW 181	std::to_string(expected_dims) + ")" +	×
NEW 182	", so " + get_name(&histogram) +	×
183	" parameter expected to be " +	×
184	std::to_string(expected_dims) +	×
185	"d. "	×
186	"Actual " +	×
NEW 187	std::to_string(histogram.get_ndim()) + "d");	×
188	}	×
189	}	616✔
190
191	int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;	756✔
192	if (histogram.get_size() > max_hist_size) {	756!
193	throw py::value_error(get_name(&histogram) +	×
194	" parameter size expected to be less than " +	×
195	std::to_string(max_hist_size) + ". Actual " +	×
196	std::to_string(histogram.get_size()));	×
197	}	×
198
199	auto array_types = dpctl_td_ns::usm_ndarray_types();	756✔
200	auto hist_type = static_cast<typenum_t>(	756✔
201	array_types.typenum_to_lookup_id(histogram.get_typenum()));	756✔
202	if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {	756!
203	auto device = exec_q.get_device();	748✔
204	bool _64bit_atomics = device.has(sycl::aspect::atomic64);	748✔
205
206	if (!_64bit_atomics) {	748!
207	auto device_name = device.get_info<sycl::info::device::name>();	×
208	throw py::value_error(	×
209	get_name(&histogram) +	×
210	" parameter has 64-bit type, but 64-bit atomics " +	×
211	" are not supported for " + device_name);	×
212	}	×
213	}	748✔
214	}	756✔
215
216	uint32_t get_local_hist_copies_count(uint32_t loc_mem_size_in_items,
217	uint32_t local_size,
218	uint32_t hist_size_in_items)
219	{	648✔
220	constexpr uint32_t local_copies_limit = 16;	648✔
221	constexpr uint32_t atomics_per_work_item = 4;	648✔
222
223	const uint32_t preferred_local_copies =	648✔
224	CeilDiv(atomics_per_work_item * local_size, hist_size_in_items);	648✔
225	const uint32_t local_copies_fit_memory =	648✔
226	loc_mem_size_in_items / hist_size_in_items;	648✔
227
228	uint32_t local_hist_count =	648✔
229	std::min(preferred_local_copies, local_copies_limit);	648✔
230	local_hist_count = std::min(local_hist_count, local_copies_fit_memory);	648✔
231
232	return local_hist_count;	648✔
233	}	648✔
234
235	} // namespace histogram
236	} // namespace statistics

IntelPython / dpnp / 11707166170

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