Skip to content

Commit 3e28657

Browse files
authored
Add functions std and var (#105)
1 parent 605416c commit 3e28657

9 files changed

+452
-19
lines changed

dpnp/backend.pxd

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -83,12 +83,14 @@ cdef extern from "backend/backend_iface_fptr.hpp" namespace "DPNPFuncName": # n
8383
DPNP_FN_SORT
8484
DPNP_FN_SQRT
8585
DPNP_FN_SQUARE
86+
DPNP_FN_STD
8687
DPNP_FN_SUBTRACT
8788
DPNP_FN_SUM
8889
DPNP_FN_TAN
8990
DPNP_FN_TANH
9091
DPNP_FN_TRANSPOSE
9192
DPNP_FN_TRUNC
93+
DPNP_FN_VAR
9294

9395
cdef extern from "backend/backend_iface_fptr.hpp" namespace "DPNPFuncType": # need this namespace for Enum import
9496
cdef enum DPNPFuncType "DPNPFuncType":

dpnp/backend/backend_iface.hpp

Lines changed: 46 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -341,6 +341,52 @@ INP_DLLEXPORT void custom_argmax_c(void* array, void* result, size_t size);
341341
template <typename _DataType, typename _idx_DataType>
342342
INP_DLLEXPORT void custom_argmin_c(void* array, void* result, size_t size);
343343

344+
/**
345+
* @ingroup BACKEND_API
346+
* @brief MKL implementation of std function
347+
*
348+
* @param [in] array Input array with data.
349+
*
350+
* @param [out] result Output array with indeces.
351+
*
352+
* @param [in] shape Shape of input array.
353+
*
354+
* @param [in] ndim Number of elements in shape.
355+
*
356+
* @param [in] axis Axis.
357+
*
358+
* @param [in] naxis Number of elements in axis.
359+
*
360+
* @param [in] ddof Delta degrees of freedom.
361+
*
362+
*/
363+
template <typename _DataType, typename _ResultType>
364+
INP_DLLEXPORT void custom_std_c(void* array, void* result, size_t* shape,
365+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
366+
367+
/**
368+
* @ingroup BACKEND_API
369+
* @brief MKL implementation of var function
370+
*
371+
* @param [in] array Input array with data.
372+
*
373+
* @param [out] result Output array with indeces.
374+
*
375+
* @param [in] shape Shape of input array.
376+
*
377+
* @param [in] ndim Number of elements in shape.
378+
*
379+
* @param [in] axis Axis.
380+
*
381+
* @param [in] naxis Number of elements in axis.
382+
*
383+
* @param [in] ddof Delta degrees of freedom.
384+
*
385+
*/
386+
template <typename _DataType, typename _ResultType>
387+
INP_DLLEXPORT void custom_var_c(void* array, void* result, size_t* shape,
388+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
389+
344390
#if 0 // Example for OpenCL kernel
345391
template <typename _DataType>
346392
void custom_dgemm_c_opencl(void* array_1, void* array_2, void* result_1, size_t size);

dpnp/backend/backend_iface_fptr.cpp

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -610,6 +610,11 @@ static func_map_t func_map_init()
610610
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_elemwise_square_c<float>};
611611
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_elemwise_square_c<double>};
612612

613+
fmap[DPNPFuncName::DPNP_FN_STD][eft_INT][eft_INT] = {eft_DBL, (void*)custom_std_c<int, double>};
614+
fmap[DPNPFuncName::DPNP_FN_STD][eft_LNG][eft_LNG] = {eft_DBL, (void*)custom_std_c<long, double>};
615+
fmap[DPNPFuncName::DPNP_FN_STD][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_std_c<float, float>};
616+
fmap[DPNPFuncName::DPNP_FN_STD][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_std_c<double, double>};
617+
613618
fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_INT] = {eft_INT,
614619
(void*)custom_elemwise_subtract_c<int, int, int>};
615620
fmap[DPNPFuncName::DPNP_FN_SUBTRACT][eft_INT][eft_LNG] = {eft_LNG,
@@ -668,5 +673,10 @@ static func_map_t func_map_init()
668673
fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_elemwise_trunc_c<float, float>};
669674
fmap[DPNPFuncName::DPNP_FN_TRUNC][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_elemwise_trunc_c<double, double>};
670675

676+
fmap[DPNPFuncName::DPNP_FN_VAR][eft_INT][eft_INT] = {eft_DBL, (void*)custom_var_c<int, double>};
677+
fmap[DPNPFuncName::DPNP_FN_VAR][eft_LNG][eft_LNG] = {eft_DBL, (void*)custom_var_c<long, double>};
678+
fmap[DPNPFuncName::DPNP_FN_VAR][eft_FLT][eft_FLT] = {eft_FLT, (void*)custom_var_c<float, float>};
679+
fmap[DPNPFuncName::DPNP_FN_VAR][eft_DBL][eft_DBL] = {eft_DBL, (void*)custom_var_c<double, double>};
680+
671681
return fmap;
672682
};

dpnp/backend/backend_iface_fptr.hpp

Lines changed: 2 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -112,12 +112,14 @@ enum class DPNPFuncName : size_t
112112
DPNP_FN_SORT, /**< Used in numpy.sort() implementation */
113113
DPNP_FN_SQRT, /**< Used in numpy.sqrt() implementation */
114114
DPNP_FN_SQUARE, /**< Used in numpy.square() implementation */
115+
DPNP_FN_STD, /**< Used in numpy.std() implementation */
115116
DPNP_FN_SUBTRACT, /**< Used in numpy.subtract() implementation */
116117
DPNP_FN_SUM, /**< Used in numpy.sum() implementation */
117118
DPNP_FN_TAN, /**< Used in numpy.tan() implementation */
118119
DPNP_FN_TANH, /**< Used in numpy.tanh() implementation */
119120
DPNP_FN_TRANSPOSE, /**< Used in numpy.transpose() implementation */
120121
DPNP_FN_TRUNC, /**< Used in numpy.trunc() implementation */
122+
DPNP_FN_VAR, /**< Used in numpy.var() implementation */
121123
DPNP_FN_LAST /**< The latest element of the enumeration */
122124
};
123125

dpnp/backend/custom_kernels_statistics.cpp

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -140,6 +140,7 @@ void custom_cov_c(void* array1_in, void* result1, size_t nrows, size_t ncols)
140140

141141
template void custom_cov_c<double>(void* array1_in, void* result1, size_t nrows, size_t ncols);
142142

143+
143144
template <typename _DataType>
144145
class custom_max_c_kernel;
145146

@@ -282,3 +283,91 @@ template void
282283
template void
283284
custom_min_c<long>(void* array1_in, void* result1, size_t* shape, size_t ndim, size_t* axis, size_t naxis);
284285
template void custom_min_c<int>(void* array1_in, void* result1, size_t* shape, size_t ndim, size_t* axis, size_t naxis);
286+
287+
288+
template <typename _DataType, typename _ResultType>
289+
void custom_std_c(void* array1_in, void* result1, size_t* shape, size_t ndim, size_t* axis, size_t naxis, size_t ddof)
290+
{
291+
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
292+
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
293+
294+
_ResultType* var = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(1 * sizeof(_ResultType)));
295+
custom_var_c<_DataType, _ResultType>(array1, var, shape, ndim, axis, naxis, ddof);
296+
297+
custom_elemwise_sqrt_c<_ResultType, _ResultType>(var, result, 1);
298+
299+
dpnp_memory_free_c(var);
300+
301+
#if 0
302+
std::cout << "std result " << res[0] << "\n";
303+
#endif
304+
}
305+
306+
template void custom_std_c<int, double>(void* array1_in, void* result1, size_t* shape,
307+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
308+
template void custom_std_c<long, double>(void* array1_in, void* result1, size_t* shape,
309+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
310+
template void custom_std_c<float, float>(void* array1_in, void* result1, size_t* shape,
311+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
312+
template void custom_std_c<double, double>(void* array1_in, void* result1, size_t* shape,
313+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
314+
315+
316+
template <typename _DataType, typename _ResultType>
317+
class custom_var_c_kernel;
318+
319+
template <typename _DataType, typename _ResultType>
320+
void custom_var_c(void* array1_in, void* result1, size_t* shape, size_t ndim, size_t* axis, size_t naxis, size_t ddof)
321+
{
322+
cl::sycl::event event;
323+
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in);
324+
_ResultType* result = reinterpret_cast<_ResultType*>(result1);
325+
326+
_ResultType* mean = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(1 * sizeof(_ResultType)));
327+
custom_mean_c<_DataType, _ResultType>(array1, mean, shape, ndim, axis, naxis);
328+
_ResultType mean_val = mean[0];
329+
330+
size_t size = 1;
331+
for (size_t i = 0; i < ndim; ++i) {
332+
size *= shape[i];
333+
}
334+
335+
_ResultType* squared_deviations = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(size * sizeof(_ResultType)));
336+
337+
cl::sycl::range<1> gws(size);
338+
event = DPNP_QUEUE.submit([&](cl::sycl::handler& cgh) {
339+
cgh.parallel_for<class custom_var_c_kernel<_DataType, _ResultType> >(
340+
gws,
341+
[=](cl::sycl::id<1> global_id)
342+
{
343+
size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/
344+
{
345+
_ResultType deviation = (_ResultType)array1[i] - mean_val;
346+
squared_deviations[i] = deviation * deviation;
347+
}
348+
}); /* parallel_for */
349+
}); /* queue.submit */
350+
351+
event.wait();
352+
353+
custom_mean_c<_ResultType, _ResultType>(squared_deviations, mean, shape, ndim, axis, naxis);
354+
mean_val = mean[0];
355+
356+
result[0] = mean_val * size / (size - ddof);
357+
358+
dpnp_memory_free_c(mean);
359+
dpnp_memory_free_c(squared_deviations);
360+
361+
#if 0
362+
std::cout << "var result " << res[0] << "\n";
363+
#endif
364+
}
365+
366+
template void custom_var_c<int, double>(void* array1_in, void* result1, size_t* shape,
367+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
368+
template void custom_var_c<long, double>(void* array1_in, void* result1, size_t* shape,
369+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
370+
template void custom_var_c<float, float>(void* array1_in, void* result1, size_t* shape,
371+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);
372+
template void custom_var_c<double, double>(void* array1_in, void* result1, size_t* shape,
373+
size_t ndim, size_t* axis, size_t naxis, size_t ddof);

dpnp/backend_statistics.pyx

Lines changed: 37 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -46,18 +46,45 @@ __all__ += [
4646
"dpnp_max",
4747
"dpnp_mean",
4848
"dpnp_median",
49-
"dpnp_min"
49+
"dpnp_min",
50+
"dpnp_std",
51+
"dpnp_var",
5052
]
5153

5254

5355
# C function pointer to the C library template functions
5456
ctypedef void(*fptr_custom_cov_1in_1out_t)(void * , void * , size_t, size_t)
55-
57+
ctypedef void(*fptr_custom_std_var_1in_1out_t)(void * , void * , size_t * , size_t, size_t * , size_t, size_t)
5658

5759
# C function pointer to the C library template functions
5860
ctypedef void(*custom_statistic_1in_1out_func_ptr_t)(void * , void * , size_t * , size_t, size_t * , size_t)
5961

6062

63+
cdef dparray call_fptr_custom_std_var_1in_1out(DPNPFuncName fptr_name, dparray a, ddof):
64+
65+
""" Convert string type names (dparray.dtype) to C enum DPNPFuncType """
66+
cdef DPNPFuncType param_type = dpnp_dtype_to_DPNPFuncType(a.dtype)
67+
68+
""" get the FPTR data structure """
69+
cdef DPNPFuncData kernel_data = get_dpnp_function_ptr(fptr_name, param_type, DPNP_FT_NONE)
70+
71+
result_type = dpnp_DPNPFuncType_to_dtype( < size_t > kernel_data.return_type)
72+
""" Create result array with type given by FPTR data """
73+
cdef dparray result = dparray((1,), dtype=result_type)
74+
75+
cdef fptr_custom_std_var_1in_1out_t func = <fptr_custom_std_var_1in_1out_t > kernel_data.ptr
76+
77+
# stub for interface support
78+
cdef dparray_shape_type axis
79+
cdef Py_ssize_t axis_size = 0
80+
81+
""" Call FPTR function """
82+
func(a.get_data(), result.get_data(), < size_t * > a._dparray_shape.data(),
83+
a.ndim, < size_t * > axis.data(), axis_size, ddof)
84+
85+
return result
86+
87+
6188
cpdef dpnp_average(dparray x1):
6289
array_sum = dpnp_sum(x1)
6390

@@ -432,3 +459,11 @@ cpdef dparray dpnp_min(dparray input, axis):
432459
dpnp_array = dpnp.array(result_array, dtype=input.dtype)
433460
dpnp_result_array = dpnp_array.reshape(output_shape)
434461
return dpnp_result_array
462+
463+
464+
cpdef dparray dpnp_std(dparray a, size_t ddof):
465+
return call_fptr_custom_std_var_1in_1out(DPNP_FN_STD, a, ddof)
466+
467+
468+
cpdef dparray dpnp_var(dparray a, size_t ddof):
469+
return call_fptr_custom_std_var_1in_1out(DPNP_FN_VAR, a, ddof)

dpnp/dparray.pyx

Lines changed: 18 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -518,6 +518,15 @@ cdef class dparray:
518518
"""
519519
return repeat(self, *args, **kwds)
520520

521+
def std(self, axis=None, dtype=None, out=None, ddof=0, keepdims=False):
522+
""" Returns the variance of the array elements, along given axis.
523+
524+
.. seealso::
525+
:func:`dpnp.var` for full documentation,
526+
527+
"""
528+
return std(self, axis, dtype, out, ddof, keepdims)
529+
521530
def transpose(self, *axes):
522531
""" Returns a view of the array with axes permuted.
523532
@@ -528,6 +537,15 @@ cdef class dparray:
528537
"""
529538
return transpose(self, axes)
530539

540+
def var(self, axis=None, dtype=None, out=None, ddof=0, keepdims=False):
541+
""" Returns the standard deviation of the array elements along given axis.
542+
543+
.. seealso::
544+
:func:`dpnp.std` for full documentation,
545+
546+
"""
547+
return var(self, axis, dtype, out, ddof, keepdims)
548+
531549
def __array_ufunc__(self, ufunc, method, *inputs, **kwargs):
532550
""" Return our function pointer regusted by `ufunc` parameter
533551
"""

0 commit comments

Comments
 (0)