12224470186

Committed 08 Dec 2024 07:04PM UTC coverage: 66.716% (-1.4%) from 68.085%

Build # 12224470186

Build Type

Pull #2180

github

Committed by

web-flow

Commit Message

Merge a90cb6fd7 into 6dc39f9e9

Pull Request Pull Request #2180: Implementation of correlate

Run Details

4487 of 10508 branches covered (42.7%)

Branch coverage included in aggregate %.

276 of 589 new or added lines in 9 files covered. (46.86%)

60 existing lines in 3 files now uncovered.

16704 of 21255 relevant lines covered (78.59%)

20105.49 hits per line

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

62.22

/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 <vector>

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

#include <pybind11/pybind11.h>

#include "histogram_common.hpp"

#include "validation_utils.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;

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

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

namespace 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);
    check_num_dims(&histogram, 1, names);

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

        auto sample_size = sample.get_size();
        auto weights_size = weights_ptr->get_size();
        if (sample.get_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);
    }
    else if (sample.get_ndim() == 2) {
        auto sample_count = sample.get_shape(0);
        auto expected_dims = sample.get_shape(1);

        if (bins_ptr != nullptr && bins_ptr->get_ndim() != expected_dims) {
            throw py::value_error(
                name_of(&sample, names) + " parameter has shape {" +
                std::to_string(sample_count) + "x" +
                std::to_string(expected_dims) + "}" + ", so " +
                name_of(bins_ptr, names) + " parameter expected to be " +
                std::to_string(expected_dims) +
                "d. "
                "Actual " +
                std::to_string(bins->get_ndim()) + "d");
        }
    }

    if (bins_ptr != nullptr) {
        py::ssize_t expected_hist_size = 1;
        for (int i = 0; i < bins_ptr->get_ndim(); ++i) {
            expected_hist_size *= (bins_ptr->get_shape(i) - 1);
        }

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

IntelPython / dpnp / 12224470186

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