Skip to content

Commit 57caae8

Browse files
shssfshssf
authored
use sycl_adapter in krnl_statistics (#904)
1 parent c7f527e commit 57caae8

3 files changed

Lines changed: 84 additions & 79 deletions

File tree

‎dpnp/backend/include/dpnp_iface.hpp‎

Lines changed: 6 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -211,10 +211,11 @@ INP_DLLEXPORT void
211211
* @param [in] array Input array.
212212
* @param [in] mask_arr Input mask array when elem is nan.
213213
* @param [out] result Output array.
214+
* @param [in] result_size Output array size.
214215
* @param [in] size Number of elements in input arrays.
215216
*/
216217
template <typename _DataType>
217-
INP_DLLEXPORT void dpnp_nanvar_c(void* array, void* mask_arr, void* result, size_t size);
218+
INP_DLLEXPORT void dpnp_nanvar_c(void* array, void* mask_arr, void* result, const size_t result_size, size_t size);
218219

219220
/**
220221
* @ingroup BACKEND_API
@@ -645,14 +646,15 @@ INP_DLLEXPORT void dpnp_matrix_rank_c(void* array1_in, void* result1, size_t* sh
645646
*
646647
* @param [in] array Input array with data.
647648
* @param [out] result Output array.
649+
* @param [in] result_size Output array size.
648650
* @param [in] shape Shape of input array.
649651
* @param [in] ndim Number of elements in shape.
650652
* @param [in] axis Axis.
651653
* @param [in] naxis Number of elements in axis.
652654
*/
653655
template <typename _DataType>
654656
INP_DLLEXPORT void
655-
dpnp_max_c(void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
657+
dpnp_max_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
656658

657659
/**
658660
* @ingroup BACKEND_API
@@ -690,14 +692,15 @@ INP_DLLEXPORT void
690692
*
691693
* @param [in] array Input array with data.
692694
* @param [out] result Output array.
695+
* @param [in] result_size Output array size.
693696
* @param [in] shape Shape of input array.
694697
* @param [in] ndim Number of elements in shape.
695698
* @param [in] axis Axis.
696699
* @param [in] naxis Number of elements in axis.
697700
*/
698701
template <typename _DataType>
699702
INP_DLLEXPORT void
700-
dpnp_min_c(void* array, void* result, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
703+
dpnp_min_c(void* array, void* result, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis);
701704

702705
/**
703706
* @ingroup BACKEND_API

‎dpnp/backend/kernels/dpnp_krnl_statistics.cpp‎

Lines changed: 65 additions & 71 deletions
Original file line numberDiff line numberDiff line change
@@ -28,6 +28,7 @@
2828
#include <dpnp_iface.hpp>
2929
#include "dpnp_fptr.hpp"
3030
#include "dpnp_utils.hpp"
31+
#include "dpnpc_memory_adapter.hpp"
3132
#include "queue_sycl.hpp"
3233

3334
namespace mkl_blas = oneapi::mkl::blas::row_major;
@@ -68,7 +69,8 @@ class dpnp_cov_c_kernel;
6869
template <typename _DataType>
6970
void dpnp_cov_c(void* array1_in, void* result1, size_t nrows, size_t ncols)
7071
{
71-
_DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
72+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, nrows * ncols);
73+
_DataType* array_1 = input1_ptr.get_ptr();
7274
_DataType* result = reinterpret_cast<_DataType*>(result1);
7375

7476
if (!nrows || !ncols)
@@ -144,15 +146,23 @@ template <typename _DataType>
144146
class dpnp_max_c_kernel;
145147

146148
template <typename _DataType>
147-
void dpnp_max_c(void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
149+
void dpnp_max_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
148150
{
151+
const size_t size_input = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
152+
if (!size_input)
153+
{
154+
return;
155+
}
156+
157+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, size_input, true);
158+
DPNPC_ptr_adapter<_DataType> result_ptr(result1, result_size, true, true);
159+
_DataType* array_1 = input1_ptr.get_ptr();
160+
_DataType* result = result_ptr.get_ptr();
161+
149162
if (naxis == 0)
150163
{
151164
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
152165

153-
_DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
154-
_DataType* result = reinterpret_cast<_DataType*>(result1);
155-
156166
size_t size = 1;
157167
for (size_t i = 0; i < ndim; ++i)
158168
{
@@ -182,9 +192,6 @@ void dpnp_max_c(void* array1_in, void* result1, const size_t* shape, size_t ndim
182192
}
183193
else
184194
{
185-
_DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
186-
_DataType* result = reinterpret_cast<_DataType*>(result1);
187-
188195
size_t res_ndim = ndim - naxis;
189196
size_t res_shape[res_ndim];
190197
int ind = 0;
@@ -206,12 +213,6 @@ void dpnp_max_c(void* array1_in, void* result1, const size_t* shape, size_t ndim
206213
}
207214
}
208215

209-
size_t size_input = 1;
210-
for (size_t i = 0; i < ndim; ++i)
211-
{
212-
size_input *= shape[i];
213-
}
214-
215216
size_t input_shape_offsets[ndim];
216217
size_t acc = 1;
217218
for (size_t i = ndim - 1; i > 0; --i)
@@ -338,24 +339,20 @@ void dpnp_mean_c(void* array1_in, void* result1, const size_t* shape, size_t ndi
338339
{
339340
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
340341

341-
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
342-
343-
size_t size = 1;
344-
for (size_t i = 0; i < ndim; ++i)
345-
{
346-
size *= shape[i];
347-
}
348-
342+
const size_t size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
349343
if (!size)
350344
{
351345
return;
352346
}
353347

348+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, size, true);
349+
DPNPC_ptr_adapter<_ResultType> result_ptr(result1, 1, true, true);
350+
_DataType* array = input1_ptr.get_ptr();
351+
_ResultType* result = result_ptr.get_ptr();
352+
354353
if constexpr (std::is_same<_DataType, double>::value || std::is_same<_DataType, float>::value)
355354
{
356-
_ResultType* array = reinterpret_cast<_DataType*>(array1_in);
357-
358-
auto dataset = mkl_stats::make_dataset<mkl_stats::layout::row_major>(1, size, array);
355+
auto dataset = mkl_stats::make_dataset<mkl_stats::layout::row_major /*, _ResultType*/>(1, size, array);
359356

360357
cl::sycl::event event = mkl_stats::mean(DPNP_QUEUE, dataset, result);
361358

@@ -366,7 +363,7 @@ void dpnp_mean_c(void* array1_in, void* result1, const size_t* shape, size_t ndi
366363
_ResultType* sum = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(1 * sizeof(_ResultType)));
367364

368365
dpnp_sum_c<_ResultType, _DataType>(
369-
sum, array1_in, shape, ndim, reinterpret_cast<const long*>(axis), naxis, nullptr, nullptr);
366+
sum, array, shape, ndim, reinterpret_cast<const long*>(axis), naxis, nullptr, nullptr);
370367

371368
result[0] = sum[0] / static_cast<_ResultType>(size);
372369

@@ -381,14 +378,15 @@ void dpnp_median_c(void* array1_in, void* result1, const size_t* shape, size_t n
381378
{
382379
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
383380

384-
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
385-
386-
size_t size = 1;
387-
for (size_t i = 0; i < ndim; ++i)
381+
const size_t size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
382+
if (!size)
388383
{
389-
size *= shape[i];
384+
return;
390385
}
391386

387+
DPNPC_ptr_adapter<_ResultType> result_ptr(result1, 1, true, true);
388+
_ResultType* result = result_ptr.get_ptr();
389+
392390
_DataType* sorted = reinterpret_cast<_DataType*>(dpnp_memory_alloc_c(size * sizeof(_DataType)));
393391

394392
dpnp_sort_c<_DataType>(array1_in, sorted, size);
@@ -411,26 +409,29 @@ template <typename _DataType>
411409
class dpnp_min_c_kernel;
412410

413411
template <typename _DataType>
414-
void dpnp_min_c(void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
412+
void dpnp_min_c(void* array1_in, void* result1, const size_t result_size, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis)
415413
{
416-
if (naxis == 0)
414+
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
415+
416+
const size_t size_input = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
417+
if (!size_input)
417418
{
418-
__attribute__((unused)) void* tmp = (void*)(axis + naxis);
419+
return;
420+
}
419421

420-
_DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
421-
_DataType* result = reinterpret_cast<_DataType*>(result1);
422+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, size_input, true);
423+
DPNPC_ptr_adapter<_DataType> result_ptr(result1, result_size, true, true);
424+
_DataType* array_1 = input1_ptr.get_ptr();
425+
_DataType* result = result_ptr.get_ptr();
422426

423-
size_t size = 1;
424-
for (size_t i = 0; i < ndim; ++i)
425-
{
426-
size *= shape[i];
427-
}
427+
if (naxis == 0)
428+
{
428429
if constexpr (std::is_same<_DataType, double>::value || std::is_same<_DataType, float>::value)
429430
{
430431
// Required initializing the result before call the function
431432
result[0] = array_1[0];
432433

433-
auto dataset = mkl_stats::make_dataset<mkl_stats::layout::row_major>(1, size, array_1);
434+
auto dataset = mkl_stats::make_dataset<mkl_stats::layout::row_major>(1, size_input, array_1);
434435

435436
cl::sycl::event event = mkl_stats::min(DPNP_QUEUE, dataset, result);
436437

@@ -440,17 +441,14 @@ void dpnp_min_c(void* array1_in, void* result1, const size_t* shape, size_t ndim
440441
{
441442
auto policy = oneapi::dpl::execution::make_device_policy<class dpnp_min_c_kernel<_DataType>>(DPNP_QUEUE);
442443

443-
_DataType* res = std::min_element(policy, array_1, array_1 + size);
444+
_DataType* res = std::min_element(policy, array_1, array_1 + size_input);
444445
policy.queue().wait();
445446

446447
result[0] = *res;
447448
}
448449
}
449450
else
450451
{
451-
_DataType* array_1 = reinterpret_cast<_DataType*>(array1_in);
452-
_DataType* result = reinterpret_cast<_DataType*>(result1);
453-
454452
size_t res_ndim = ndim - naxis;
455453
size_t res_shape[res_ndim];
456454
int ind = 0;
@@ -472,12 +470,6 @@ void dpnp_min_c(void* array1_in, void* result1, const size_t* shape, size_t ndim
472470
}
473471
}
474472

475-
size_t size_input = 1;
476-
for (size_t i = 0; i < ndim; ++i)
477-
{
478-
size_input *= shape[i];
479-
}
480-
481473
size_t input_shape_offsets[ndim];
482474
size_t acc = 1;
483475
for (size_t i = ndim - 1; i > 0; --i)
@@ -600,13 +592,9 @@ void dpnp_min_c(void* array1_in, void* result1, const size_t* shape, size_t ndim
600592
}
601593

602594
template <typename _DataType>
603-
void dpnp_nanvar_c(void* array1_in, void* mask_arr1, void* result1, size_t arr_size)
595+
void dpnp_nanvar_c(void* array1_in, void* mask_arr1, void* result1, const size_t result_size, size_t arr_size)
604596
{
605-
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
606-
bool* mask_arr = reinterpret_cast<bool*>(mask_arr1);
607-
_DataType* result = reinterpret_cast<_DataType*>(result1);
608-
609-
if ((array1 == nullptr) || (mask_arr == nullptr) || (result == nullptr))
597+
if ((array1_in == nullptr) || (mask_arr1 == nullptr) || (result1 == nullptr))
610598
{
611599
return;
612600
}
@@ -616,6 +604,13 @@ void dpnp_nanvar_c(void* array1_in, void* mask_arr1, void* result1, size_t arr_s
616604
return;
617605
}
618606

607+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, arr_size, true);
608+
DPNPC_ptr_adapter<bool> input2_ptr(mask_arr1, arr_size, true);
609+
DPNPC_ptr_adapter<_DataType> result_ptr(result1, result_size, true, true);
610+
_DataType* array1 = input1_ptr.get_ptr();
611+
bool* mask_arr = input2_ptr.get_ptr();
612+
_DataType* result = result_ptr.get_ptr();
613+
619614
size_t ind = 0;
620615
for (size_t i = 0; i < arr_size; ++i)
621616
{
@@ -633,13 +628,10 @@ template <typename _DataType, typename _ResultType>
633628
void dpnp_std_c(
634629
void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis, size_t ddof)
635630
{
636-
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
637-
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
638-
639631
_ResultType* var = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(1 * sizeof(_ResultType)));
640-
dpnp_var_c<_DataType, _ResultType>(array1, var, shape, ndim, axis, naxis, ddof);
641632

642-
dpnp_sqrt_c<_ResultType, _ResultType>(var, result, 1);
633+
dpnp_var_c<_DataType, _ResultType>(array1_in, var, shape, ndim, axis, naxis, ddof);
634+
dpnp_sqrt_c<_ResultType, _ResultType>(var, result1, 1);
643635

644636
dpnp_memory_free_c(var);
645637

@@ -653,20 +645,22 @@ template <typename _DataType, typename _ResultType>
653645
void dpnp_var_c(
654646
void* array1_in, void* result1, const size_t* shape, size_t ndim, const size_t* axis, size_t naxis, size_t ddof)
655647
{
648+
const size_t size = std::accumulate(shape, shape + ndim, 1, std::multiplies<size_t>());
649+
if (!size)
650+
{
651+
return;
652+
}
653+
656654
cl::sycl::event event;
657-
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
658-
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
655+
DPNPC_ptr_adapter<_DataType> input1_ptr(array1_in, size);
656+
DPNPC_ptr_adapter<_ResultType> result_ptr(result1, 1, true, true);
657+
_DataType* array1 = input1_ptr.get_ptr();
658+
_ResultType* result = result_ptr.get_ptr();
659659

660660
_ResultType* mean = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(1 * sizeof(_ResultType)));
661661
dpnp_mean_c<_DataType, _ResultType>(array1, mean, shape, ndim, axis, naxis);
662662
_ResultType mean_val = mean[0];
663663

664-
size_t size = 1;
665-
for (size_t i = 0; i < ndim; ++i)
666-
{
667-
size *= shape[i];
668-
}
669-
670664
_ResultType* squared_deviations = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(size * sizeof(_ResultType)));
671665

672666
cl::sycl::range<1> gws(size);

‎dpnp/dpnp_algo/dpnp_algo_statistics.pyx‎

Lines changed: 13 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -50,11 +50,12 @@ __all__ += [
5050

5151
# C function pointer to the C library template functions
5252
ctypedef void(*fptr_custom_cov_1in_1out_t)(void *, void * , size_t, size_t)
53-
ctypedef void(*fptr_custom_nanvar_t)(void *, void * , void * , size_t)
53+
ctypedef void(*fptr_custom_nanvar_t)(void *, void * , void * , size_t, size_t)
5454
ctypedef void(*fptr_custom_std_var_1in_1out_t)(void *, void * , size_t * , size_t, size_t * , size_t, size_t)
5555

5656
# C function pointer to the C library template functions
5757
ctypedef void(*custom_statistic_1in_1out_func_ptr_t)(void *, void * , size_t * , size_t, size_t * , size_t)
58+
ctypedef void(*custom_statistic_1in_1out_func_ptr_t_max)(void *, void * , const size_t, size_t * , size_t, size_t * , size_t)
5859

5960

6061
cdef utils.dpnp_descriptor call_fptr_custom_std_var_1in_1out(DPNPFuncName fptr_name, utils.dpnp_descriptor x1, ddof):
@@ -146,7 +147,7 @@ cdef utils.dpnp_descriptor _dpnp_max(utils.dpnp_descriptor input, _axis_, shape_
146147
# ceate result array with type given by FPTR data
147148
cdef utils.dpnp_descriptor result = utils.create_output_descriptor(result_shape, kernel_data.return_type, None)
148149

149-
cdef custom_statistic_1in_1out_func_ptr_t func = <custom_statistic_1in_1out_func_ptr_t > kernel_data.ptr
150+
cdef custom_statistic_1in_1out_func_ptr_t_max func = <custom_statistic_1in_1out_func_ptr_t_max > kernel_data.ptr
150151
cdef shape_type_c axis
151152
cdef Py_ssize_t axis_size = 0
152153
cdef shape_type_c axis_ = axis
@@ -158,7 +159,13 @@ cdef utils.dpnp_descriptor _dpnp_max(utils.dpnp_descriptor input, _axis_, shape_
158159
axis_.push_back(shape_it)
159160
axis_size = len(axis)
160161

161-
func(input.get_data(), result.get_data(), < size_t * > input_shape.data(), input.ndim, < size_t * > axis_.data(), axis_size)
162+
func(input.get_data(),
163+
result.get_data(),
164+
result.size,
165+
< size_t * > input_shape.data(),
166+
input.ndim,
167+
< size_t * > axis_.data(),
168+
axis_size)
162169

163170
return result
164171

@@ -355,7 +362,7 @@ cpdef utils.dpnp_descriptor _dpnp_min(utils.dpnp_descriptor input, _axis_, shape
355362

356363
cdef utils.dpnp_descriptor result = utils.create_output_descriptor(shape_output, kernel_data.return_type, None)
357364

358-
cdef custom_statistic_1in_1out_func_ptr_t func = <custom_statistic_1in_1out_func_ptr_t > kernel_data.ptr
365+
cdef custom_statistic_1in_1out_func_ptr_t_max func = <custom_statistic_1in_1out_func_ptr_t_max > kernel_data.ptr
359366
cdef shape_type_c axis
360367
cdef Py_ssize_t axis_size = 0
361368
cdef shape_type_c axis_ = axis
@@ -371,6 +378,7 @@ cpdef utils.dpnp_descriptor _dpnp_min(utils.dpnp_descriptor input, _axis_, shape
371378

372379
func(input.get_data(),
373380
result.get_data(),
381+
result.size,
374382
< size_t * > input_shape.data(),
375383
input.ndim,
376384
< size_t * > axis_.data(),
@@ -423,7 +431,7 @@ cpdef utils.dpnp_descriptor dpnp_nanvar(utils.dpnp_descriptor arr, ddof):
423431

424432
cdef fptr_custom_nanvar_t func = <fptr_custom_nanvar_t > kernel_data.ptr
425433

426-
func(arr.get_data(), mask_arr.get_data(), result.get_data(), arr.size)
434+
func(arr.get_data(), mask_arr.get_data(), result.get_data(), result.size, arr.size)
427435

428436
return call_fptr_custom_std_var_1in_1out(DPNP_FN_VAR, result, ddof)
429437

0 commit comments

Comments
 (0)