11247863672

Committed 09 Oct 2024 03:46AM UTC coverage: 57.304% (-2.2%) from 59.534%

Build # 11247863672

Build Type

Pull #2027

github

Committed by

web-flow

Commit Message

Merge ad8291ff6 into 48515c86f

Pull Request Pull Request #2027: Implementation of histogram with sycl kernel

Run Details

3867 of 10710 branches covered (36.11%)

Branch coverage included in aggregate %.

260 of 497 new or added lines in 5 files covered. (52.31%)

1 existing line in 1 file now uncovered.

14036 of 20532 relevant lines covered (68.36%)

16415.55 hits per line

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

46.05

/dpnp/backend/extensions/sycl_ext/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/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 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 + "'";
    };

    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");
    }

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

    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_size();
        auto weights_size = weights_ptr->get_size();
        if (sample.get_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 (bins.get_ndim() != 1) {
            throw py::value_error(get_name(&sample) + " parameter is 1d, but " +
                                  get_name(&bins) + " is " +
                                  std::to_string(bins.get_ndim()) + "d");
        }
    }
    else if (sample.get_ndim() == 2) {
        auto sample_count = sample.get_shape(0);
        auto expected_dims = sample.get_shape(1);

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

    py::ssize_t expected_hist_size = 1;
    for (int i = 0; i < bins.get_ndim(); ++i) {
        expected_hist_size *= (bins.get_shape(i) - 1);
    }

    if (histogram.get_size() != expected_hist_size) {
        throw py::value_error(
            get_name(&histogram) + " and " + get_name(&bins) +
            " shape mismatch. " + get_name(&histogram) +
            " 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(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);
        }
    }
}

} // namespace 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 <limits>
28	#include <string>
29	#include <unordered_map>
30	#include <vector>
31
32	#include "dpctl4pybind11.hpp"
33	#include "utils/type_dispatch.hpp"
34	#include <pybind11/pybind11.h>
35
36	#include "histogram_common.hpp"
37
38	namespace dpctl_td_ns = dpctl::tensor::type_dispatch;
39	using dpctl::tensor::usm_ndarray;
40	using dpctl_td_ns::typenum_t;
41
42	namespace histogram
43	{
44
45	void validate(const usm_ndarray &sample,
46	const usm_ndarray &bins,
47	std::optional<const dpctl::tensor::usm_ndarray> &weights,
48	const usm_ndarray &histogram)
49	{	140✔
50	auto exec_q = sample.get_queue();	140✔
51	using array_ptr = const usm_ndarray *;	140✔
52
53	std::vector<array_ptr> arrays{&sample, &bins, &histogram};	140✔
54	std::unordered_map<array_ptr, std::string> names = {	140✔
55	{arrays[0], "sample"}, {arrays[1], "bins"}, {arrays[2], "histogram"}};	140✔
56
57	array_ptr weights_ptr = nullptr;	140✔
58
59	if (weights.has_value()) {	140✔
60	weights_ptr = &weights.value();	51✔
61	arrays.push_back(weights_ptr);	51✔
62	names.insert({weights_ptr, "weights"});	51✔
63	}	51✔
64
65	auto get_name = [&](const array_ptr &arr) {	140✔
NEW 66	auto name_it = names.find(arr);	×
NEW 67	assert(name_it != names.end());	×
68
NEW 69	return "'" + name_it->second + "'";	×
NEW 70	};	×
71
72	auto unequal_queue =	140✔
73	std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {	471✔
74	return arr->get_queue() != exec_q;	471✔
75	});	471✔
76
77	if (unequal_queue != arrays.cend()) {	140!
NEW 78	throw py::value_error(	×
NEW 79	get_name(*unequal_queue) +	×
NEW 80	" parameter has incompatible queue with parameter " +	×
NEW 81	get_name(&sample));	×
NEW 82	}	×
83
84	auto non_contig_array =	140✔
85	std::find_if(arrays.cbegin(), arrays.cend(), [&](const array_ptr &arr) {	471✔
86	return !arr->is_c_contiguous();	471✔
87	});	471✔
88
89	if (non_contig_array != arrays.cend()) {	140!
NEW 90	throw py::value_error(get_name(*non_contig_array) +	×
NEW 91	" parameter is not c-contiguos");	×
NEW 92	}	×
93
94	if (histogram.get_ndim() != 1) {	140!
NEW 95	throw py::value_error(get_name(&histogram) +	×
NEW 96	" parameter must be 1d. Actual " +	×
NEW 97	std::to_string(histogram.get_ndim()) + "d");	×
NEW 98	}	×
99
100	if (weights_ptr) {	140✔
101	if (weights_ptr->get_ndim() != 1) {	51!
NEW 102	throw py::value_error(	×
NEW 103	get_name(weights_ptr) + " parameter must be 1d. Actual " +	×
NEW 104	std::to_string(weights_ptr->get_ndim()) + "d");	×
NEW 105	}	×
106
107	auto sample_size = sample.get_size();	51✔
108	auto weights_size = weights_ptr->get_size();	51✔
109	if (sample.get_size() != weights_ptr->get_size()) {	51!
NEW 110	throw py::value_error(	×
NEW 111	get_name(&sample) + " size (" + std::to_string(sample_size) +	×
NEW 112	") and " + get_name(weights_ptr) + " size (" +	×
NEW 113	std::to_string(weights_size) + ")" + " must match");	×
NEW 114	}	×
115	}	51✔
116
117	if (sample.get_ndim() > 2) {	140!
NEW 118	throw py::value_error(	×
NEW 119	get_name(&sample) +	×
NEW 120	" parameter must have no more than 2 dimensions. Actual " +	×
NEW 121	std::to_string(sample.get_ndim()) + "d");	×
NEW 122	}	×
123
124	if (sample.get_ndim() == 1) {	140!
125	if (bins.get_ndim() != 1) {	140!
NEW 126	throw py::value_error(get_name(&sample) + " parameter is 1d, but " +	×
NEW 127	get_name(&bins) + " is " +	×
NEW 128	std::to_string(bins.get_ndim()) + "d");	×
NEW 129	}	×
130	}	140✔
NEW 131	else if (sample.get_ndim() == 2) {	×
NEW 132	auto sample_count = sample.get_shape(0);	×
NEW 133	auto expected_dims = sample.get_shape(1);	×
134
NEW 135	if (bins.get_ndim() != expected_dims) {	×
NEW 136	throw py::value_error(get_name(&sample) + " parameter has shape {" +	×
NEW 137	std::to_string(sample_count) + "x" +	×
NEW 138	std::to_string(expected_dims) + "}" +	×
NEW 139	", so " + get_name(&bins) +	×
NEW 140	" parameter expected to be " +	×
NEW 141	std::to_string(expected_dims) +	×
NEW 142	"d. "	×
NEW 143	"Actual " +	×
NEW 144	std::to_string(bins.get_ndim()) + "d");	×
NEW 145	}	×
NEW 146	}	×
147
148	py::ssize_t expected_hist_size = 1;	140✔
149	for (int i = 0; i < bins.get_ndim(); ++i) {	280✔
150	expected_hist_size *= (bins.get_shape(i) - 1);	140✔
151	}	140✔
152
153	if (histogram.get_size() != expected_hist_size) {	140!
NEW 154	throw py::value_error(	×
NEW 155	get_name(&histogram) + " and " + get_name(&bins) +	×
NEW 156	" shape mismatch. " + get_name(&histogram) +	×
NEW 157	" expected to have size = " + std::to_string(expected_hist_size) +	×
NEW 158	". Actual " + std::to_string(histogram.get_size()));	×
NEW 159	}	×
160
161	int64_t max_hist_size = std::numeric_limits<uint32_t>::max() - 1;	140✔
162	if (histogram.get_size() > max_hist_size) {	140!
NEW 163	throw py::value_error(get_name(&histogram) +	×
NEW 164	" parameter size expected to be less than " +	×
NEW 165	std::to_string(max_hist_size) + ". Actual " +	×
NEW 166	std::to_string(histogram.get_size()));	×
NEW 167	}	×
168
169	auto array_types = dpctl_td_ns::usm_ndarray_types();	140✔
170	auto hist_type = static_cast<typenum_t>(	140✔
171	array_types.typenum_to_lookup_id(histogram.get_typenum()));	140✔
172	if (histogram.get_elemsize() == 8 && hist_type != typenum_t::CFLOAT) {	140!
173	auto device = exec_q.get_device();	130✔
174	bool _64bit_atomics = device.has(sycl::aspect::atomic64);	130✔
175
176	if (!_64bit_atomics) {	130!
NEW 177	auto device_name = device.get_info<sycl::info::device::name>();	×
NEW 178	throw py::value_error(	×
NEW 179	get_name(&histogram) +	×
NEW 180	" parameter has 64-bit type, but 64-bit atomics " +	×
NEW 181	" are not supported for " + device_name);	×
NEW 182	}	×
183	}	130✔
184	}	140✔
185
186	} // namespace histogram

IntelPython / dpnp / 11247863672

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