14219089871

Committed 02 Apr 2025 12:02PM UTC coverage: 72.271% (+0.04%) from 72.236%

Build # 14219089871

Build Type

push

github

Committed by

web-flow

Commit Message

Disable w/a for dpctl-2030 on CPU device (#2405)

The PR proposes to partly merge changes from #2403 to unblock coverage
GitHub action.
The workaround for dpctl-2030 was kept only for testing with GPU device
(the use case from internal CI, where the changes from dpctl-2030 is not
available).

Run Details

4718 of 9534 branches covered (49.49%)

Branch coverage included in aggregate %.

17728 of 21524 relevant lines covered (82.36%)

19868.85 hits per line

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

63.51

/dpnp/backend/extensions/statistics/histogram_common.cpp

//*****************************************************************************
// Copyright (c) 2024-2025, 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 <vector>

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

#include <pybind11/pybind11.h>

#include "histogram_common.hpp"

#include "ext/validation_utils.hpp"

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

using ext::common::CeilDiv;

using ext::validation::array_names;
using ext::validation::array_ptr;

using ext::validation::check_max_dims;
using ext::validation::check_num_dims;
using ext::validation::check_size_at_least;
using ext::validation::common_checks;
using ext::validation::name_of;

namespace statistics::histogram
{

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

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

    array_ptr bins_ptr = nullptr;

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

    array_ptr weights_ptr = nullptr;

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

    common_checks({&sample, bins.has_value() ? &bins.value() : nullptr,
                   weights.has_value() ? &weights.value() : nullptr},
                  {&histogram}, names);

    check_size_at_least(bins_ptr, 2, names);
    check_size_at_least(&histogram, 1, names);

    if (weights_ptr) {
        check_num_dims(weights_ptr, 1, names);

        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(name_of(&sample, names) + " size (" +
                                  std::to_string(sample_size) + ") and " +
                                  name_of(weights_ptr, names) + " size (" +
                                  std::to_string(weights_size) + ")" +
                                  " must match");
        }
    }

    check_max_dims(&sample, 2, names);

    if (sample.get_ndim() == 1) {
        check_num_dims(bins_ptr, 1, names);

        if (bins_ptr && histogram.get_size() != bins_ptr->get_size() - 1) {
            auto hist_size = histogram.get_size();
            auto bins_size = bins_ptr->get_size();
            throw py::value_error(
                name_of(&histogram, names) + " parameter and " +
                name_of(bins_ptr, names) + " parameters shape mismatch. " +
                name_of(&histogram, names) + " size is " +
                std::to_string(hist_size) + name_of(bins_ptr, names) +
                " 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(
                name_of(&sample, names) + " parameter has shape (" +
                std::to_string(sample_count) + ", " +
                std::to_string(expected_dims) + ")" + ", so " +
                name_of(&histogram, names) + " parameter expected to be " +
                std::to_string(expected_dims) +
                "d. "
                "Actual " +
                std::to_string(histogram.get_ndim()) + "d");
        }

        if (bins_ptr != nullptr) {
            py::ssize_t expected_bins_size = 0;
            for (int i = 0; i < histogram.get_ndim(); ++i) {
                expected_bins_size += histogram.get_shape(i) + 1;
            }

            auto actual_bins_size = bins_ptr->get_size();
            if (actual_bins_size != expected_bins_size) {
                throw py::value_error(
                    name_of(&histogram, names) + " and " +
                    name_of(bins_ptr, names) + " shape mismatch. " +
                    name_of(bins_ptr, names) + " expected to have size = " +
                    std::to_string(expected_bins_size) + ". Actual " +
                    std::to_string(actual_bins_size));
            }
        }

        int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;
        if (histogram.get_size() > max_hist_size) {
            throw py::value_error(name_of(&histogram, names) +
                                  " 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(
                    name_of(&histogram, names) +
                    " 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 statistics::histogram

1	//*****************************************************************************
2	// Copyright (c) 2024-2025, 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 <vector>
30
31	#include "dpctl4pybind11.hpp"
32	#include "utils/type_dispatch.hpp"
33
34	#include <pybind11/pybind11.h>
35
36	#include "histogram_common.hpp"
37
38	#include "ext/validation_utils.hpp"
39
40	namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
41	using dpctl::tensor::usm_ndarray;
42	using dpctl_td_ns::typenum_t;
43
44	using ext::common::CeilDiv;
45
46	using ext::validation::array_names;
47	using ext::validation::array_ptr;
48
49	using ext::validation::check_max_dims;
50	using ext::validation::check_num_dims;
51	using ext::validation::check_size_at_least;
52	using ext::validation::common_checks;
53	using ext::validation::name_of;
54
55	namespace statistics::histogram
56	{
57
58	void validate(const usm_ndarray &sample,
59	const std::optional<const dpctl::tensor::usm_ndarray> &bins,
60	const std::optional<const dpctl::tensor::usm_ndarray> &weights,
61	const usm_ndarray &histogram)
62	{	1,387✔
63	auto exec_q = sample.get_queue();	1,387✔
64
65	std::vector<array_ptr> arrays{&sample, &histogram};	1,387✔
66	array_names names = {{arrays[0], "sample"}, {arrays[1], "histogram"}};	1,387✔
67
68	array_ptr bins_ptr = nullptr;	1,387✔
69
70	if (bins.has_value()) {	1,387✔
71	bins_ptr = &bins.value();	1,285✔
72	arrays.push_back(bins_ptr);	1,285✔
73	names.insert({bins_ptr, "bins"});	1,285✔
74	}	1,285✔
75
76	array_ptr weights_ptr = nullptr;	1,387✔
77
78	if (weights.has_value()) {	1,387✔
79	weights_ptr = &weights.value();	827✔
80	arrays.push_back(weights_ptr);	827✔
81	names.insert({weights_ptr, "weights"});	827✔
82	}	827✔
83
84	common_checks({&sample, bins.has_value() ? &bins.value() : nullptr,	1,387✔
85	weights.has_value() ? &weights.value() : nullptr},	1,387✔
86	{&histogram}, names);	1,387✔
87
88	check_size_at_least(bins_ptr, 2, names);	1,387✔
89	check_size_at_least(&histogram, 1, names);	1,387✔
90
91	if (weights_ptr) {	1,387✔
92	check_num_dims(weights_ptr, 1, names);	827✔
93
94	auto sample_size = sample.get_shape(0);	827✔
95	auto weights_size = weights_ptr->get_size();	827✔
96	if (sample_size != weights_ptr->get_size()) {	827!
97	throw py::value_error(name_of(&sample, names) + " size (" +	×
98	std::to_string(sample_size) + ") and " +	×
99	name_of(weights_ptr, names) + " size (" +	×
100	std::to_string(weights_size) + ")" +	×
101	" must match");	×
102	}	×
103	}	827✔
104
105	check_max_dims(&sample, 2, names);	1,387✔
106
107	if (sample.get_ndim() == 1) {	1,387✔
108	check_num_dims(bins_ptr, 1, names);	238✔
109
110	if (bins_ptr && histogram.get_size() != bins_ptr->get_size() - 1) {	238!
111	auto hist_size = histogram.get_size();	×
112	auto bins_size = bins_ptr->get_size();	×
113	throw py::value_error(	×
114	name_of(&histogram, names) + " parameter and " +	×
115	name_of(bins_ptr, names) + " parameters shape mismatch. " +	×
116	name_of(&histogram, names) + " size is " +	×
117	std::to_string(hist_size) + name_of(bins_ptr, names) +	×
118	" must have size " + std::to_string(hist_size + 1) +	×
119	" but have " + std::to_string(bins_size));	×
120	}	×
121	}	238✔
122	else if (sample.get_ndim() == 2) {	1,149!
123	auto sample_count = sample.get_shape(0);	1,149✔
124	auto expected_dims = sample.get_shape(1);	1,149✔
125
126	if (histogram.get_ndim() != expected_dims) {	1,149!
127	throw py::value_error(	×
128	name_of(&sample, names) + " parameter has shape (" +	×
129	std::to_string(sample_count) + ", " +	×
130	std::to_string(expected_dims) + ")" + ", so " +	×
131	name_of(&histogram, names) + " parameter expected to be " +	×
132	std::to_string(expected_dims) +	×
133	"d. "	×
134	"Actual " +	×
135	std::to_string(histogram.get_ndim()) + "d");	×
136	}	×
137
138	if (bins_ptr != nullptr) {	1,149!
139	py::ssize_t expected_bins_size = 0;	1,149✔
140	for (int i = 0; i < histogram.get_ndim(); ++i) {	4,017✔
141	expected_bins_size += histogram.get_shape(i) + 1;	2,868✔
142	}	2,868✔
143
144	auto actual_bins_size = bins_ptr->get_size();	1,149✔
145	if (actual_bins_size != expected_bins_size) {	1,149!
146	throw py::value_error(	×
147	name_of(&histogram, names) + " and " +	×
148	name_of(bins_ptr, names) + " shape mismatch. " +	×
149	name_of(bins_ptr, names) + " expected to have size = " +	×
150	std::to_string(expected_bins_size) + ". Actual " +	×
151	std::to_string(actual_bins_size));	×
152	}	×
153	}	1,149✔
154
155	int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;	1,149✔
156	if (histogram.get_size() > max_hist_size) {	1,149!
157	throw py::value_error(name_of(&histogram, names) +	×
158	" parameter size expected to be less than " +	×
159	std::to_string(max_hist_size) + ". Actual " +	×
160	std::to_string(histogram.get_size()));	×
161	}	×
162
163	auto array_types = dpctl_td_ns::usm_ndarray_types();	1,149✔
164	auto hist_type = static_cast<typenum_t>(	1,149✔
165	array_types.typenum_to_lookup_id(histogram.get_typenum()));	1,149✔
166	if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {	1,149!
167	auto device = exec_q.get_device();	1,149✔
168	bool _64bit_atomics = device.has(sycl::aspect::atomic64);	1,149✔
169
170	if (!_64bit_atomics) {	1,149!
171	auto device_name = device.get_info<sycl::info::device::name>();	×
172	throw py::value_error(	×
173	name_of(&histogram, names) +	×
174	" parameter has 64-bit type, but 64-bit atomics " +	×
175	" are not supported for " + device_name);	×
176	}	×
177	}	1,149✔
178	}	1,149✔
179	}	1,387✔
180
181	uint32_t get_local_hist_copies_count(uint32_t loc_mem_size_in_items,
182	uint32_t local_size,
183	uint32_t hist_size_in_items)
184	{	1,373✔
185	constexpr uint32_t local_copies_limit = 16;	1,373✔
186	constexpr uint32_t atomics_per_work_item = 4;	1,373✔
187
188	const uint32_t preferred_local_copies =	1,373✔
189	CeilDiv(atomics_per_work_item * local_size, hist_size_in_items);	1,373✔
190	const uint32_t local_copies_fit_memory =	1,373✔
191	loc_mem_size_in_items / hist_size_in_items;	1,373✔
192
193	uint32_t local_hist_count =	1,373✔
194	std::min(preferred_local_copies, local_copies_limit);	1,373✔
195	local_hist_count = std::min(local_hist_count, local_copies_fit_memory);	1,373✔
196
197	return local_hist_count;	1,373✔
198	}	1,373✔
199
200	} // namespace statistics::histogram

IntelPython / dpnp / 14219089871

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