10937375650

Committed 19 Sep 2024 08:33AM UTC coverage: 59.467% (+0.02%) from 59.452%

Build # 10937375650

Build Type

Pull #2063

github

Committed by

web-flow

Commit Message

Merge be09527fd into 40e4052a0

Pull Request Pull Request #2063: Add a protection thorugh extra memory copy in backend implementation of `dpnp.choose`

Run Details

3746 of 9432 branches covered (39.72%)

Branch coverage included in aggregate %.

3 of 3 new or added lines in 1 file covered. (100.0%)

6 existing lines in 1 file now uncovered.

13757 of 20001 relevant lines covered (68.78%)

16848.57 hits per line

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

50.0

/dpnp/backend/kernels/dpnp_krnl_indexing.cpp

//*****************************************************************************
// Copyright (c) 2016-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 <iostream>
#include <list>
#include <vector>

#include "dpnp_fptr.hpp"
#include "dpnpc_memory_adapter.hpp"
#include "queue_sycl.hpp"
#include <dpnp_iface.hpp>

template <typename _DataType1, typename _DataType2>
class dpnp_choose_c_kernel;

template <typename _DataType1, typename _DataType2>
DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref,
                                void *result1,
                                void *array1_in,
                                void **choices1,
                                size_t size,
                                size_t choices_size,
                                size_t choice_size,
                                const DPCTLEventVectorRef dep_event_vec_ref)
{
    // avoid warning unused variable
    (void)dep_event_vec_ref;

    DPCTLSyclEventRef event_ref = nullptr;

    if ((array1_in == nullptr) || (result1 == nullptr) || (choices1 == nullptr))
    {
        return event_ref;
    }
    if (!size || !choices_size || !choice_size) {
        return event_ref;
    }

    sycl::queue q = *(reinterpret_cast<sycl::queue *>(q_ref));

    DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size);
    _DataType1 *array_in = input1_ptr.get_ptr();

    DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size);
    _DataType2 **choices = choices_ptr.get_ptr();

    for (size_t i = 0; i < choices_size; ++i) {
        DPNPC_ptr_adapter<_DataType2> choice_ptr(q_ref, choices[i],
                                                 choice_size);
        choices[i] = choice_ptr.get_ptr();
    }

    DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, true,
                                              true);
    _DataType2 *result = result1_ptr.get_ptr();

    sycl::range<1> gws(size);
    auto kernel_parallel_for_func = [=](sycl::id<1> global_id) {
        const size_t idx = global_id[0];
        result[idx] = choices[array_in[idx]][idx];
    };

    auto kernel_func = [&](sycl::handler &cgh) {
        cgh.parallel_for<class dpnp_choose_c_kernel<_DataType1, _DataType2>>(
            gws, kernel_parallel_for_func);
    };

    sycl::event event = q.submit(kernel_func);
    result1_ptr.depends_on(event);

    event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);

    return DPCTLEvent_Copy(event_ref);
}

template <typename _DataType1, typename _DataType2>
void dpnp_choose_c(void *result1,
                   void *array1_in,
                   void **choices1,
                   size_t size,
                   size_t choices_size,
                   size_t choice_size)
{
    DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);
    DPCTLEventVectorRef dep_event_vec_ref = nullptr;
    DPCTLSyclEventRef event_ref = dpnp_choose_c<_DataType1, _DataType2>(
        q_ref, result1, array1_in, choices1, size, choices_size, choice_size,
        dep_event_vec_ref);
    DPCTLEvent_WaitAndThrow(event_ref);
}

template <typename _DataType1, typename _DataType2>
void (*dpnp_choose_default_c)(void *, void *, void **, size_t, size_t, size_t) =
    dpnp_choose_c<_DataType1, _DataType2>;

template <typename _DataType1, typename _DataType2>
DPCTLSyclEventRef (*dpnp_choose_ext_c)(DPCTLSyclQueueRef,
                                       void *,
                                       void *,
                                       void **,
                                       size_t,
                                       size_t,
                                       size_t,
                                       const DPCTLEventVectorRef) =
    dpnp_choose_c<_DataType1, _DataType2>;

void func_map_init_indexing_func(func_map_t &fmap)
{
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_INT] = {
        eft_INT, (void *)dpnp_choose_default_c<int32_t, int32_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_LNG] = {
        eft_LNG, (void *)dpnp_choose_default_c<int32_t, int64_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_FLT] = {
        eft_FLT, (void *)dpnp_choose_default_c<int32_t, float>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_DBL] = {
        eft_DBL, (void *)dpnp_choose_default_c<int32_t, double>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_INT] = {
        eft_INT, (void *)dpnp_choose_default_c<int64_t, int32_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_LNG] = {
        eft_LNG, (void *)dpnp_choose_default_c<int64_t, int64_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_FLT] = {
        eft_FLT, (void *)dpnp_choose_default_c<int64_t, float>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_DBL] = {
        eft_DBL, (void *)dpnp_choose_default_c<int64_t, double>};

    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_INT] = {
        eft_INT, (void *)dpnp_choose_ext_c<int32_t, int32_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_LNG] = {
        eft_LNG, (void *)dpnp_choose_ext_c<int32_t, int64_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_FLT] = {
        eft_FLT, (void *)dpnp_choose_ext_c<int32_t, float>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_DBL] = {
        eft_DBL, (void *)dpnp_choose_ext_c<int32_t, double>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_INT] = {
        eft_INT, (void *)dpnp_choose_ext_c<int64_t, int32_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_LNG] = {
        eft_LNG, (void *)dpnp_choose_ext_c<int64_t, int64_t>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_FLT] = {
        eft_FLT, (void *)dpnp_choose_ext_c<int64_t, float>};
    fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_DBL] = {
        eft_DBL, (void *)dpnp_choose_ext_c<int64_t, double>};
    return;
}

1	//*****************************************************************************
2	// Copyright (c) 2016-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 <iostream>
27	#include <list>
28	#include <vector>
29
30	#include "dpnp_fptr.hpp"
31	#include "dpnpc_memory_adapter.hpp"
32	#include "queue_sycl.hpp"
33	#include <dpnp_iface.hpp>
34
35	template <typename _DataType1, typename _DataType2>
36	class dpnp_choose_c_kernel;
37
38	template <typename _DataType1, typename _DataType2>
39	DPCTLSyclEventRef dpnp_choose_c(DPCTLSyclQueueRef q_ref,
40	void *result1,
41	void *array1_in,
42	void **choices1,
43	size_t size,
44	size_t choices_size,
45	size_t choice_size,
46	const DPCTLEventVectorRef dep_event_vec_ref)
47	{	4✔
48	// avoid warning unused variable
49	(void)dep_event_vec_ref;	4✔
50
51	DPCTLSyclEventRef event_ref = nullptr;	4✔
52
53	if ((array1_in == nullptr) \|\| (result1 == nullptr) \|\| (choices1 == nullptr))	4!
54	{	×
55	return event_ref;	×
56	}	×
57	if (!size \|\| !choices_size \|\| !choice_size) {	4!
58	return event_ref;	×
59	}	×
60
61	sycl::queue q = (reinterpret_cast<sycl::queue >(q_ref));	4✔
62
63	DPNPC_ptr_adapter<_DataType1> input1_ptr(q_ref, array1_in, size);	4✔
64	_DataType1 *array_in = input1_ptr.get_ptr();	4✔
65
66	DPNPC_ptr_adapter<_DataType2 *> choices_ptr(q_ref, choices1, choices_size);	4✔
67	_DataType2 **choices = choices_ptr.get_ptr();	4✔
68
69	for (size_t i = 0; i < choices_size; ++i) {	16!
70	DPNPC_ptr_adapter<_DataType2> choice_ptr(q_ref, choices[i],	12✔
71	choice_size);	12✔
72	choices[i] = choice_ptr.get_ptr();	12✔
73	}	12✔
74
75	DPNPC_ptr_adapter<_DataType2> result1_ptr(q_ref, result1, size, true,	4✔
76	true);	4✔
77	_DataType2 *result = result1_ptr.get_ptr();	4✔
78
79	sycl::range<1> gws(size);	4✔
80	auto kernel_parallel_for_func = [=](sycl::id<1> global_id) {	4✔
UNCOV 81	const size_t idx = global_id[0];	×
UNCOV 82	result[idx] = choices[array_in[idx]][idx];	×
UNCOV 83	};	×
84
85	auto kernel_func = [&](sycl::handler &cgh) {	4✔
86	cgh.parallel_for<class dpnp_choose_c_kernel<_DataType1, _DataType2>>(	4✔
87	gws, kernel_parallel_for_func);	4✔
88	};	4✔
89
90	sycl::event event = q.submit(kernel_func);	4✔
91	result1_ptr.depends_on(event);	4✔
92
93	event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);	4✔
94
95	return DPCTLEvent_Copy(event_ref);	4✔
96	}	4✔
97
98	template <typename _DataType1, typename _DataType2>
99	void dpnp_choose_c(void *result1,
100	void *array1_in,
101	void **choices1,
102	size_t size,
103	size_t choices_size,
104	size_t choice_size)
UNCOV 105	{	×
UNCOV 106	DPCTLSyclQueueRef q_ref = reinterpret_cast<DPCTLSyclQueueRef>(&DPNP_QUEUE);	×
UNCOV 107	DPCTLEventVectorRef dep_event_vec_ref = nullptr;	×
108	DPCTLSyclEventRef event_ref = dpnp_choose_c<_DataType1, _DataType2>(	×
109	q_ref, result1, array1_in, choices1, size, choices_size, choice_size,	×
110	dep_event_vec_ref);	×
111	DPCTLEvent_WaitAndThrow(event_ref);	×
112	}	×
113
114	template <typename _DataType1, typename _DataType2>
115	void (dpnp_choose_default_c)(void , void , void *, size_t, size_t, size_t) =
116	dpnp_choose_c<_DataType1, _DataType2>;
117
118	template <typename _DataType1, typename _DataType2>
119	DPCTLSyclEventRef (*dpnp_choose_ext_c)(DPCTLSyclQueueRef,
120	void *,
121	void *,
122	void **,
123	size_t,
124	size_t,
125	size_t,
126	const DPCTLEventVectorRef) =
127	dpnp_choose_c<_DataType1, _DataType2>;
128
129	void func_map_init_indexing_func(func_map_t &fmap)
130	{	1✔
131	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_INT] = {	1✔
132	eft_INT, (void *)dpnp_choose_default_c<int32_t, int32_t>};	1✔
133	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_LNG] = {	1✔
134	eft_LNG, (void *)dpnp_choose_default_c<int32_t, int64_t>};	1✔
135	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_FLT] = {	1✔
136	eft_FLT, (void *)dpnp_choose_default_c<int32_t, float>};	1✔
137	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_INT][eft_DBL] = {	1✔
138	eft_DBL, (void *)dpnp_choose_default_c<int32_t, double>};	1✔
139	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_INT] = {	1✔
140	eft_INT, (void *)dpnp_choose_default_c<int64_t, int32_t>};	1✔
141	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_LNG] = {	1✔
142	eft_LNG, (void *)dpnp_choose_default_c<int64_t, int64_t>};	1✔
143	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_FLT] = {	1✔
144	eft_FLT, (void *)dpnp_choose_default_c<int64_t, float>};	1✔
145	fmap[DPNPFuncName::DPNP_FN_CHOOSE][eft_LNG][eft_DBL] = {	1✔
146	eft_DBL, (void *)dpnp_choose_default_c<int64_t, double>};	1✔
147
148	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_INT] = {	1✔
149	eft_INT, (void *)dpnp_choose_ext_c<int32_t, int32_t>};	1✔
150	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_LNG] = {	1✔
151	eft_LNG, (void *)dpnp_choose_ext_c<int32_t, int64_t>};	1✔
152	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_FLT] = {	1✔
153	eft_FLT, (void *)dpnp_choose_ext_c<int32_t, float>};	1✔
154	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_INT][eft_DBL] = {	1✔
155	eft_DBL, (void *)dpnp_choose_ext_c<int32_t, double>};	1✔
156	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_INT] = {	1✔
157	eft_INT, (void *)dpnp_choose_ext_c<int64_t, int32_t>};	1✔
158	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_LNG] = {	1✔
159	eft_LNG, (void *)dpnp_choose_ext_c<int64_t, int64_t>};	1✔
160	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_FLT] = {	1✔
161	eft_FLT, (void *)dpnp_choose_ext_c<int64_t, float>};	1✔
162	fmap[DPNPFuncName::DPNP_FN_CHOOSE_EXT][eft_LNG][eft_DBL] = {	1✔
163	eft_DBL, (void *)dpnp_choose_ext_c<int64_t, double>};	1✔
164	return;	1✔
165	}	1✔

IntelPython / dpnp / 10937375650

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