Skip to content

Fix take #1167

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 4 commits into from
Aug 18, 2022
Merged

Fix take #1167

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
4 commits
Select commit Hold shift + click to select a range
885f91a
dpnp_take_c uses SYCL kernel, no need to use no_sycl parameter in ada…
oleksandr-pavlyk Aug 11, 2022
7220d87
Change to DPNPC_adapter to set/use events upon which deallocation mus…
oleksandr-pavlyk Aug 11, 2022
37386bb
Used DPNPC_ptr_adapter::depends_on
oleksandr-pavlyk Aug 14, 2022
d839ea1
implemented PR feedback
oleksandr-pavlyk Aug 17, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
47 changes: 29 additions & 18 deletions dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -493,6 +493,9 @@ DPCTLSyclEventRef dpnp_ptp_c(DPCTLSyclQueueRef q_ref,
(void)dep_event_vec_ref;

DPCTLSyclEventRef event_ref = nullptr;
DPCTLSyclEventRef e1_ref = nullptr;
DPCTLSyclEventRef e2_ref = nullptr;
DPCTLSyclEventRef e3_ref = nullptr;

if ((input1_in == nullptr) || (result1_out == nullptr))
{
Expand All @@ -514,29 +517,36 @@ DPCTLSyclEventRef dpnp_ptp_c(DPCTLSyclQueueRef q_ref,
_DataType* min_arr = reinterpret_cast<_DataType*>(sycl::malloc_shared(result_size * sizeof(_DataType), q));
_DataType* max_arr = reinterpret_cast<_DataType*>(sycl::malloc_shared(result_size * sizeof(_DataType), q));

dpnp_min_c<_DataType>(arr, min_arr, result_size, input_shape, input_ndim, axis, naxis);
dpnp_max_c<_DataType>(arr, max_arr, result_size, input_shape, input_ndim, axis, naxis);
e1_ref = dpnp_min_c<_DataType>(q_ref, arr, min_arr, result_size, input_shape, input_ndim, axis, naxis, NULL);
e2_ref = dpnp_max_c<_DataType>(q_ref, arr, max_arr, result_size, input_shape, input_ndim, axis, naxis, NULL);

shape_elem_type* _strides =
reinterpret_cast<shape_elem_type*>(sycl::malloc_shared(result_ndim * sizeof(shape_elem_type), q));
get_shape_offsets_inkernel(result_shape, result_ndim, _strides);

dpnp_subtract_c<_DataType, _DataType, _DataType>(result,
result_size,
result_ndim,
result_shape,
result_strides,
max_arr,
result_size,
result_ndim,
result_shape,
_strides,
min_arr,
result_size,
result_ndim,
result_shape,
_strides,
NULL);
e3_ref = dpnp_subtract_c<_DataType, _DataType, _DataType>(q_ref, result,
result_size,
result_ndim,
result_shape,
result_strides,
max_arr,
result_size,
result_ndim,
result_shape,
_strides,
min_arr,
result_size,
result_ndim,
result_shape,
_strides,
NULL, NULL);

DPCTLEvent_Wait(e1_ref);
DPCTLEvent_Wait(e2_ref);
DPCTLEvent_Wait(e3_ref);
DPCTLEvent_Delete(e1_ref);
DPCTLEvent_Delete(e2_ref);
DPCTLEvent_Delete(e3_ref);

sycl::free(min_arr, q);
sycl::free(max_arr, q);
Expand Down Expand Up @@ -576,6 +586,7 @@ void dpnp_ptp_c(void* result1_out,
naxis,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType>
Expand Down
24 changes: 16 additions & 8 deletions dpnp/backend/kernels/dpnp_krnl_bitwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -148,16 +148,16 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
\
sycl::queue q = *(reinterpret_cast<sycl::queue*>(q_ref)); \
\
DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, input1_in, input1_size); \
DPNPC_ptr_adapter<shape_elem_type> input1_shape_ptr(q_ref, input1_shape, input1_ndim, true); \
DPNPC_ptr_adapter<shape_elem_type> input1_strides_ptr(q_ref, input1_strides, input1_ndim, true); \
DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, input1_in, input1_size); \
DPNPC_ptr_adapter<shape_elem_type> input1_shape_ptr(q_ref, input1_shape, input1_ndim, true); \
DPNPC_ptr_adapter<shape_elem_type> input1_strides_ptr(q_ref, input1_strides, input1_ndim, true); \
\
DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, input2_in, input2_size); \
DPNPC_ptr_adapter<shape_elem_type> input2_shape_ptr(q_ref, input2_shape, input2_ndim, true); \
DPNPC_ptr_adapter<shape_elem_type> input2_strides_ptr(q_ref, input2_strides, input2_ndim, true); \
DPNPC_ptr_adapter<_DataType> input2_ptr(q_ref, input2_in, input2_size); \
DPNPC_ptr_adapter<shape_elem_type> input2_shape_ptr(q_ref, input2_shape, input2_ndim, true); \
DPNPC_ptr_adapter<shape_elem_type> input2_strides_ptr(q_ref, input2_strides, input2_ndim, true); \
\
DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result_out, result_size, false, true); \
DPNPC_ptr_adapter<shape_elem_type> result_strides_ptr(q_ref, result_strides, result_ndim); \
DPNPC_ptr_adapter<_DataType> result_ptr(q_ref, result_out, result_size, false, true); \
DPNPC_ptr_adapter<shape_elem_type> result_strides_ptr(q_ref, result_strides, result_ndim); \
\
_DataType* input1_data = input1_ptr.get_ptr(); \
shape_elem_type* input1_shape_data = input1_shape_ptr.get_ptr(); \
Expand Down Expand Up @@ -226,6 +226,14 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
}; \
event = q.submit(kernel_func); \
} \
input1_ptr.depends_on(event); \
input1_shape_ptr.depends_on(event); \
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Looks both input1_shape_ptr and input2_shape_ptr aren't used in the kernel. Should the dependency on event be created for them anyway?

input1_strides_ptr.depends_on(event); \
input2_ptr.depends_on(event); \
input2_shape_ptr.depends_on(event); \
input2_strides_ptr.depends_on(event); \
result_ptr.depends_on(event); \
result_strides_ptr.depends_on(event); \
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
\
return DPCTLEvent_Copy(event_ref); \
Expand Down
23 changes: 23 additions & 0 deletions dpnp/backend/kernels/dpnp_krnl_elemwise.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -143,6 +143,12 @@
} \
} \
\
input1_ptr.depends_on(event); \
input1_shape_ptr.depends_on(event); \
input1_strides_ptr.depends_on(event); \
result_ptr.depends_on(event); \
result_strides_ptr.depends_on(event); \
\
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
\
return DPCTLEvent_Copy(event_ref); \
Expand Down Expand Up @@ -644,6 +650,12 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
} \
} \
\
input1_ptr.depends_on(event); \
input1_shape_ptr.depends_on(event); \
input1_strides_ptr.depends_on(event); \
result_ptr.depends_on(event); \
result_strides_ptr.depends_on(event); \
\
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
\
return DPCTLEvent_Copy(event_ref); \
Expand Down Expand Up @@ -998,6 +1010,17 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
event = q.submit(kernel_func); \
} \
} \
\
input1_ptr.depends_on(event); \
input1_shape_ptr.depends_on(event); \
input1_strides_ptr.depends_on(event); \
input2_ptr.depends_on(event); \
input2_shape_ptr.depends_on(event); \
input2_strides_ptr.depends_on(event); \
result_ptr.depends_on(event); \
result_shape_ptr.depends_on(event); \
result_strides_ptr.depends_on(event); \
\
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
\
return DPCTLEvent_Copy(event_ref); \
Expand Down
2 changes: 1 addition & 1 deletion dpnp/backend/kernels/dpnp_krnl_indexing.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -901,7 +901,7 @@ DPCTLSyclEventRef dpnp_take_c(DPCTLSyclQueueRef q_ref,
DPCTLSyclEventRef event_ref = nullptr;
sycl::queue q = *(reinterpret_cast<sycl::queue*>(q_ref));

DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, array1_size, true);
DPNPC_ptr_adapter<_DataType> input1_ptr(q_ref, array1_in, array1_size);
DPNPC_ptr_adapter<_IndecesType> input2_ptr(q_ref, indices1, size);
_DataType* array_1 = input1_ptr.get_ptr();
_IndecesType* indices = input2_ptr.get_ptr();
Expand Down
8 changes: 8 additions & 0 deletions dpnp/backend/kernels/dpnp_krnl_mathematical.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -170,6 +170,8 @@ DPCTLSyclEventRef dpnp_elemwise_absolute_c(DPCTLSyclQueueRef q_ref,
event = q.submit(kernel_func);
}

input1_ptr.depends_on(event);
result1_ptr.depends_on(event);
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);

return DPCTLEvent_Copy(event_ref);
Expand Down Expand Up @@ -483,6 +485,8 @@ DPCTLSyclEventRef dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
};
event = q.submit(kernel_func);

input1_ptr.depends_on(event);
result_ptr.depends_on(event);
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);

return DPCTLEvent_Copy(event_ref);
Expand Down Expand Up @@ -676,6 +680,7 @@ void dpnp_floor_divide_c(void* result_out,
where,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_output, typename _DataType_input1, typename _DataType_input2>
Expand Down Expand Up @@ -770,6 +775,7 @@ void dpnp_modf_c(void* array1_in, void* result1_out, void* result2_out, size_t s
size,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_input, typename _DataType_output>
Expand Down Expand Up @@ -911,6 +917,7 @@ void dpnp_remainder_c(void* result_out,
where,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_output, typename _DataType_input1, typename _DataType_input2>
Expand Down Expand Up @@ -1041,6 +1048,7 @@ void dpnp_trapz_c(
array2_size,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_input1, typename _DataType_input2, typename _DataType_output>
Expand Down
2 changes: 2 additions & 0 deletions dpnp/backend/kernels/dpnp_krnl_reduction.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -162,6 +162,7 @@ void dpnp_sum_c(void* result_out,
where,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_output, typename _DataType_input>
Expand Down Expand Up @@ -278,6 +279,7 @@ void dpnp_prod_c(void* result_out,
where,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType_output, typename _DataType_input>
Expand Down
4 changes: 4 additions & 0 deletions dpnp/backend/kernels/dpnp_krnl_sorting.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -91,6 +91,7 @@ void dpnp_argsort_c(void* array1_in, void* result1, size_t size)
size,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType, typename _idx_DataType>
Expand Down Expand Up @@ -242,6 +243,7 @@ void dpnp_partition_c(
ndim,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType>
Expand Down Expand Up @@ -394,6 +396,7 @@ void dpnp_searchsorted_c(
v_size,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType, typename _IndexingType>
Expand Down Expand Up @@ -459,6 +462,7 @@ void dpnp_sort_c(void* array1_in, void* result1, size_t size)
size,
dep_event_vec_ref);
DPCTLEvent_WaitAndThrow(event_ref);
DPCTLEvent_Delete(event_ref);
}

template <typename _DataType>
Expand Down
Loading