12895869998

Committed 21 Jan 2025 09:07PM UTC coverage: 71.211% (+0.4%) from 70.856%

Build # 12895869998

Build Type

Pull #2201

github

Committed by

web-flow

Commit Message

Merge cf50357ca into 356184a29

Pull Request Pull Request #2201: Implement extension for `dpnp.choose`

Run Details

4568 of 9390 branches covered (48.65%)

Branch coverage included in aggregate %.

282 of 333 new or added lines in 5 files covered. (84.68%)

4 existing lines in 3 files now uncovered.

16935 of 20806 relevant lines covered (81.39%)

20542.24 hits per line

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

61.22

/dpnp/backend/extensions/indexing/choose_kernel.hpp

//*****************************************************************************
// 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.
//*****************************************************************************

#pragma once
#include <algorithm>
#include <complex>
#include <cstdint>
#include <limits>
#include <sycl/sycl.hpp>
#include <type_traits>

#include "kernels/dpctl_tensor_types.hpp"
#include "utils/indexing_utils.hpp"
#include "utils/offset_utils.hpp"
#include "utils/strided_iters.hpp"
#include "utils/type_utils.hpp"

namespace dpnp::extensions::indexing::strides_detail
{

struct NthStrideOffsetUnpacked
{
    NthStrideOffsetUnpacked(int common_nd,
                            dpctl::tensor::ssize_t const *_offsets,
                            dpctl::tensor::ssize_t const *_shape,
                            dpctl::tensor::ssize_t const *_strides)
        : _ind(common_nd), nd(common_nd), offsets(_offsets), shape(_shape),
          strides(_strides)
    {
    }

    template <typename nT>
    size_t operator()(dpctl::tensor::ssize_t gid, nT n) const
    {
        dpctl::tensor::ssize_t relative_offset(0);
        _ind.get_displacement<const dpctl::tensor::ssize_t *,
                              const dpctl::tensor::ssize_t *>(
            gid, shape, strides + (n * nd), relative_offset);

        return relative_offset + offsets[n];
    }

private:
    dpctl::tensor::strides::CIndexer_vector<dpctl::tensor::ssize_t> _ind;

    int nd;
    dpctl::tensor::ssize_t const *offsets;
    dpctl::tensor::ssize_t const *shape;
    dpctl::tensor::ssize_t const *strides;
};

} // namespace dpnp::extensions::indexing::strides_detail

namespace dpnp::extensions::indexing::kernels
{

template <typename ProjectorT,
          typename IndOutIndexerT,
          typename ChoicesIndexerT,
          typename IndT,
          typename T>
class ChooseFunctor
{
private:
    const IndT *ind = nullptr;
    T *dst = nullptr;
    char **chcs = nullptr;
    dpctl::tensor::ssize_t n_chcs;
    const IndOutIndexerT ind_out_indexer;
    const ChoicesIndexerT chcs_indexer;

public:
    ChooseFunctor(const IndT *ind_,
                  T *dst_,
                  char **chcs_,
                  dpctl::tensor::ssize_t n_chcs_,
                  const IndOutIndexerT &ind_out_indexer_,
                  const ChoicesIndexerT &chcs_indexer_)
        : ind(ind_), dst(dst_), chcs(chcs_), n_chcs(n_chcs_),
          ind_out_indexer(ind_out_indexer_), chcs_indexer(chcs_indexer_)
    {
    }

    void operator()(sycl::id<1> id) const
    {
        const ProjectorT proj{};

        dpctl::tensor::ssize_t i = id[0];

        auto ind_dst_offsets = ind_out_indexer(i);
        dpctl::tensor::ssize_t ind_offset = ind_dst_offsets.get_first_offset();
        dpctl::tensor::ssize_t dst_offset = ind_dst_offsets.get_second_offset();

        IndT chc_idx = ind[ind_offset];
        // proj produces an index in the range of n_chcs
        dpctl::tensor::ssize_t projected_idx = proj(n_chcs, chc_idx);

        dpctl::tensor::ssize_t chc_offset = chcs_indexer(i, projected_idx);

        T *chc = reinterpret_cast<T *>(chcs[projected_idx]);

        dst[dst_offset] = chc[chc_offset];
    }
};

typedef sycl::event (*choose_fn_ptr_t)(sycl::queue &,
                                       size_t,
                                       dpctl::tensor::ssize_t,
                                       int,
                                       const dpctl::tensor::ssize_t *,
                                       const char *,
                                       char *,
                                       char **,
                                       dpctl::tensor::ssize_t,
                                       dpctl::tensor::ssize_t,
                                       const dpctl::tensor::ssize_t *,
                                       const std::vector<sycl::event> &);

template <typename ProjectorT, typename indTy, typename Ty>
sycl::event choose_impl(sycl::queue &q,
                        size_t nelems,
                        dpctl::tensor::ssize_t n_chcs,
                        int nd,
                        const dpctl::tensor::ssize_t *shape_and_strides,
                        const char *ind_cp,
                        char *dst_cp,
                        char **chcs_cp,
                        dpctl::tensor::ssize_t ind_offset,
                        dpctl::tensor::ssize_t dst_offset,
                        const dpctl::tensor::ssize_t *chc_offsets,
                        const std::vector<sycl::event> &depends)
{
    dpctl::tensor::type_utils::validate_type_for_device<Ty>(q);

    const indTy *ind_tp = reinterpret_cast<const indTy *>(ind_cp);
    Ty *dst_tp = reinterpret_cast<Ty *>(dst_cp);

    sycl::event choose_ev = q.submit([&](sycl::handler &cgh) {
        cgh.depends_on(depends);

        using InOutIndexerT =
            dpctl::tensor::offset_utils::TwoOffsets_StridedIndexer;
        const InOutIndexerT ind_out_indexer{nd, ind_offset, dst_offset,
                                            shape_and_strides};

        using NthChoiceIndexerT = strides_detail::NthStrideOffsetUnpacked;
        const NthChoiceIndexerT choices_indexer{
            nd, chc_offsets, shape_and_strides, shape_and_strides + 3 * nd};

        using ChooseFunc = ChooseFunctor<ProjectorT, InOutIndexerT,
                                         NthChoiceIndexerT, indTy, Ty>;

        cgh.parallel_for<ChooseFunc>(sycl::range<1>(nelems),
                                     ChooseFunc(ind_tp, dst_tp, chcs_cp, n_chcs,
                                                ind_out_indexer,
                                                choices_indexer));
    });

    return choose_ev;
}

} // namespace dpnp::extensions::indexing::kernels

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	#pragma once
27	#include <algorithm>
28	#include <complex>
29	#include <cstdint>
30	#include <limits>
31	#include <sycl/sycl.hpp>
32	#include <type_traits>
33
34	#include "kernels/dpctl_tensor_types.hpp"
35	#include "utils/indexing_utils.hpp"
36	#include "utils/offset_utils.hpp"
37	#include "utils/strided_iters.hpp"
38	#include "utils/type_utils.hpp"
39
40	namespace dpnp::extensions::indexing::strides_detail
41	{
42
43	struct NthStrideOffsetUnpacked
44	{
45	NthStrideOffsetUnpacked(int common_nd,
46	dpctl::tensor::ssize_t const *_offsets,
47	dpctl::tensor::ssize_t const *_shape,
48	dpctl::tensor::ssize_t const *_strides)
49	: _ind(common_nd), nd(common_nd), offsets(_offsets), shape(_shape),	93✔
50	strides(_strides)	93✔
51	{	93✔
52	}	93✔
53
54	template <typename nT>
55	size_t operator()(dpctl::tensor::ssize_t gid, nT n) const
NEW 56	{	×
NEW 57	dpctl::tensor::ssize_t relative_offset(0);	×
NEW 58	_ind.get_displacement<const dpctl::tensor::ssize_t *,	×
NEW 59	const dpctl::tensor::ssize_t *>(	×
NEW 60	gid, shape, strides + (n * nd), relative_offset);	×
61
NEW 62	return relative_offset + offsets[n];	×
NEW 63	}	×
64
65	private:
66	dpctl::tensor::strides::CIndexer_vector<dpctl::tensor::ssize_t> _ind;
67
68	int nd;
69	dpctl::tensor::ssize_t const *offsets;
70	dpctl::tensor::ssize_t const *shape;
71	dpctl::tensor::ssize_t const *strides;
72	};
73
74	} // namespace dpnp::extensions::indexing::strides_detail
75
76	namespace dpnp::extensions::indexing::kernels
77	{
78
79	template <typename ProjectorT,
80	typename IndOutIndexerT,
81	typename ChoicesIndexerT,
82	typename IndT,
83	typename T>
84	class ChooseFunctor
85	{
86	private:
87	const IndT *ind = nullptr;
88	T *dst = nullptr;
89	char **chcs = nullptr;
90	dpctl::tensor::ssize_t n_chcs;
91	const IndOutIndexerT ind_out_indexer;
92	const ChoicesIndexerT chcs_indexer;
93
94	public:
95	ChooseFunctor(const IndT *ind_,
96	T *dst_,
97	char **chcs_,
98	dpctl::tensor::ssize_t n_chcs_,
99	const IndOutIndexerT &ind_out_indexer_,
100	const ChoicesIndexerT &chcs_indexer_)
101	: ind(ind_), dst(dst_), chcs(chcs_), n_chcs(n_chcs_),	93✔
102	ind_out_indexer(ind_out_indexer_), chcs_indexer(chcs_indexer_)	93✔
103	{	93✔
104	}	93✔
105
106	void operator()(sycl::id<1> id) const
NEW 107	{	×
NEW 108	const ProjectorT proj{};	×
109
NEW 110	dpctl::tensor::ssize_t i = id[0];	×
111
NEW 112	auto ind_dst_offsets = ind_out_indexer(i);	×
NEW 113	dpctl::tensor::ssize_t ind_offset = ind_dst_offsets.get_first_offset();	×
NEW 114	dpctl::tensor::ssize_t dst_offset = ind_dst_offsets.get_second_offset();	×
115
NEW 116	IndT chc_idx = ind[ind_offset];	×
117	// proj produces an index in the range of n_chcs
NEW 118	dpctl::tensor::ssize_t projected_idx = proj(n_chcs, chc_idx);	×
119
NEW 120	dpctl::tensor::ssize_t chc_offset = chcs_indexer(i, projected_idx);	×
121
NEW 122	T chc = reinterpret_cast<T >(chcs[projected_idx]);	×
123
NEW 124	dst[dst_offset] = chc[chc_offset];	×
NEW 125	}	×
126	};
127
128	typedef sycl::event (*choose_fn_ptr_t)(sycl::queue &,
129	size_t,
130	dpctl::tensor::ssize_t,
131	int,
132	const dpctl::tensor::ssize_t *,
133	const char *,
134	char *,
135	char **,
136	dpctl::tensor::ssize_t,
137	dpctl::tensor::ssize_t,
138	const dpctl::tensor::ssize_t *,
139	const std::vector<sycl::event> &);
140
141	template <typename ProjectorT, typename indTy, typename Ty>
142	sycl::event choose_impl(sycl::queue &q,
143	size_t nelems,
144	dpctl::tensor::ssize_t n_chcs,
145	int nd,
146	const dpctl::tensor::ssize_t *shape_and_strides,
147	const char *ind_cp,
148	char *dst_cp,
149	char **chcs_cp,
150	dpctl::tensor::ssize_t ind_offset,
151	dpctl::tensor::ssize_t dst_offset,
152	const dpctl::tensor::ssize_t *chc_offsets,
153	const std::vector<sycl::event> &depends)
154	{	93✔
155	dpctl::tensor::type_utils::validate_type_for_device<Ty>(q);	93✔
156
157	const indTy ind_tp = reinterpret_cast<const indTy >(ind_cp);	93✔
158	Ty dst_tp = reinterpret_cast<Ty >(dst_cp);	93✔
159
160	sycl::event choose_ev = q.submit([&](sycl::handler &cgh) {	93✔
161	cgh.depends_on(depends);	93✔
162
163	using InOutIndexerT =	93✔
164	dpctl::tensor::offset_utils::TwoOffsets_StridedIndexer;	93✔
165	const InOutIndexerT ind_out_indexer{nd, ind_offset, dst_offset,	93✔
166	shape_and_strides};	93✔
167
168	using NthChoiceIndexerT = strides_detail::NthStrideOffsetUnpacked;	93✔
169	const NthChoiceIndexerT choices_indexer{	93✔
170	nd, chc_offsets, shape_and_strides, shape_and_strides + 3 * nd};	93✔
171
172	using ChooseFunc = ChooseFunctor<ProjectorT, InOutIndexerT,	93✔
173	NthChoiceIndexerT, indTy, Ty>;	93✔
174
175	cgh.parallel_for<ChooseFunc>(sycl::range<1>(nelems),	93✔
176	ChooseFunc(ind_tp, dst_tp, chcs_cp, n_chcs,	93✔
177	ind_out_indexer,	93✔
178	choices_indexer));	93✔
179	});	93✔
180
181	return choose_ev;	93✔
182	}	93✔
183
184	} // namespace dpnp::extensions::indexing::kernels

IntelPython / dpnp / 12895869998

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