Skip to content

Commit 9859554

Browse files
Alexander-MakaryevAlexander-Makaryev
authored
Kernel copy func (#539)
* kernel for copy func
1 parent b6d0c4a commit 9859554

5 files changed

Lines changed: 30 additions & 11 deletions

File tree

‎dpnp/backend/include/dpnp_iface_fptr.hpp‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -78,6 +78,7 @@ enum class DPNPFuncName : size_t
7878
DPNP_FN_CEIL, /**< Used in numpy.ceil() implementation */
7979
DPNP_FN_CHOLESKY, /**< Used in numpy.linalg.cholesky() implementation */
8080
DPNP_FN_CONJIGUATE, /**< Used in numpy.conjugate() implementation */
81+
DPNP_FN_COPY, /**< Used in numpy.copy() implementation */
8182
DPNP_FN_COPYSIGN, /**< Used in numpy.copysign() implementation */
8283
DPNP_FN_CORRELATE, /**< Used in numpy.correlate() implementation */
8384
DPNP_FN_COS, /**< Used in numpy.cos() implementation */

‎dpnp/backend/kernels/dpnp_krnl_elemwise.cpp‎

Lines changed: 22 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -127,15 +127,15 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
127127
fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_degrees_c<float, float>};
128128
fmap[DPNPFuncName::DPNP_FN_DEGREES][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_degrees_c<double, double>};
129129

130-
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_exp2_c<double, double>};
131-
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_exp2_c<float, float>};
132130
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_INT][eft_INT] = {eft_DBL, (void*)dpnp_exp2_c<int, double>};
133131
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_LNG][eft_LNG] = {eft_DBL, (void*)dpnp_exp2_c<long, double>};
132+
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_exp2_c<float, float>};
133+
fmap[DPNPFuncName::DPNP_FN_EXP2][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_exp2_c<double, double>};
134134

135-
fmap[DPNPFuncName::DPNP_FN_EXP][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_exp_c<double, double>};
136-
fmap[DPNPFuncName::DPNP_FN_EXP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_exp_c<float, float>};
137135
fmap[DPNPFuncName::DPNP_FN_EXP][eft_INT][eft_INT] = {eft_DBL, (void*)dpnp_exp_c<int, double>};
138136
fmap[DPNPFuncName::DPNP_FN_EXP][eft_LNG][eft_LNG] = {eft_DBL, (void*)dpnp_exp_c<long, double>};
137+
fmap[DPNPFuncName::DPNP_FN_EXP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_exp_c<float, float>};
138+
fmap[DPNPFuncName::DPNP_FN_EXP][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_exp_c<double, double>};
139139

140140
fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_INT][eft_INT] = {eft_DBL, (void*)dpnp_expm1_c<int, double>};
141141
fmap[DPNPFuncName::DPNP_FN_EXPM1][eft_LNG][eft_LNG] = {eft_DBL, (void*)dpnp_expm1_c<long, double>};
@@ -219,6 +219,11 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
219219
{ \
220220
cl::sycl::event event; \
221221
\
222+
if (!size) \
223+
{ \
224+
return; \
225+
} \
226+
\
222227
_DataType* array1 = reinterpret_cast<_DataType*>(array1_in); \
223228
_DataType* result = reinterpret_cast<_DataType*>(result1); \
224229
\
@@ -297,25 +302,32 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
297302
fmap[DPNPFuncName::DPNP_FN_CONJIGUATE][eft_C128][eft_C128] = {eft_C128,
298303
(void*)dpnp_conjugate_c<std::complex<double>>};
299304

305+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_BOOL][eft_BOOL] = {eft_BOOL, (void*)dpnp_copy_c<bool>};
306+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_copy_c<int>};
307+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_copy_c<long>};
308+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_copy_c<float>};
309+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_copy_c<double>};
310+
fmap[DPNPFuncName::DPNP_FN_COPY][eft_C128][eft_C128] = {eft_C128, (void*)dpnp_copy_c<std::complex<double>>};
311+
300312
fmap[DPNPFuncName::DPNP_FN_ERF][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_erf_c<int>};
301313
fmap[DPNPFuncName::DPNP_FN_ERF][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_erf_c<long>};
302314
fmap[DPNPFuncName::DPNP_FN_ERF][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_erf_c<float>};
303315
fmap[DPNPFuncName::DPNP_FN_ERF][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_erf_c<double>};
304316

305-
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_recip_c<double>};
306-
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_recip_c<float>};
307317
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_recip_c<int>};
308318
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_recip_c<long>};
319+
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_recip_c<float>};
320+
fmap[DPNPFuncName::DPNP_FN_RECIP][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_recip_c<double>};
309321

310-
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_sign_c<double>};
311-
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_sign_c<float>};
312322
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_sign_c<int>};
313323
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_sign_c<long>};
324+
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_sign_c<float>};
325+
fmap[DPNPFuncName::DPNP_FN_SIGN][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_sign_c<double>};
314326

315-
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_square_c<double>};
316-
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_square_c<float>};
317327
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_square_c<int>};
318328
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_square_c<long>};
329+
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_FLT][eft_FLT] = {eft_FLT, (void*)dpnp_square_c<float>};
330+
fmap[DPNPFuncName::DPNP_FN_SQUARE][eft_DBL][eft_DBL] = {eft_DBL, (void*)dpnp_square_c<double>};
319331

320332
return;
321333
}

‎dpnp/dpnp_algo/dpnp_algo.pxd‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -51,6 +51,7 @@ cdef extern from "dpnp_iface_fptr.hpp" namespace "DPNPFuncName": # need this na
5151
DPNP_FN_CEIL
5252
DPNP_FN_CHOLESKY
5353
DPNP_FN_CONJIGUATE
54+
DPNP_FN_COPY
5455
DPNP_FN_COPYSIGN
5556
DPNP_FN_CORRELATE
5657
DPNP_FN_COS

‎dpnp/dpnp_algo/dpnp_algo_arraycreation.pyx‎

Lines changed: 5 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -41,6 +41,7 @@ from dpnp.dpnp_algo cimport *
4141

4242

4343
__all__ += [
44+
"dpnp_copy",
4445
"dpnp_diag",
4546
"dpnp_geomspace",
4647
"dpnp_linspace",
@@ -52,6 +53,10 @@ __all__ += [
5253
]
5354

5455

56+
cpdef dparray dpnp_copy(dparray x1, order, subok):
57+
return call_fptr_1in_1out(DPNP_FN_COPY, x1, x1.shape)
58+
59+
5560
cpdef dparray dpnp_diag(v, k):
5661
cdef dparray result
5762

‎dpnp/dpnp_iface_arraycreation.py‎

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -350,7 +350,7 @@ def copy(a, order='C', subok=False):
350350
"""
351351

352352
if not use_origin_backend(a):
353-
return array(a, order=order, subok=subok)
353+
return dpnp_copy(a, order, subok)
354354

355355
return call_origin(numpy.copy, a, order, subok)
356356

0 commit comments

Comments
 (0)