Skip to content

Commit c9f977f

Browse files
densmirnoleksandr-pavlyk
densmirn
authored and
oleksandr-pavlyk
committed
Extand main cython funcs with queues and events
1 parent 3411f28 commit c9f977f

17 files changed

+736
-238
lines changed

‎dpnp/backend/include/dpnp_iface_fptr.hpp

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -381,7 +381,6 @@ enum class DPNPFuncName : size_t
381381
DPNP_FN_ZEROS_LIKE, /**< Used in numpy.zeros_like() impl */
382382
DPNP_FN_ZEROS_LIKE_EXT, /**< Used in numpy.zeros_like() impl, requires extra parameters */
383383
DPNP_FN_LAST, /**< The latest element of the enumeration */
384-
DPNP_FN_LAST_EXT /**< The latest element of the enumeration, requires extra parameters */
385384
};
386385

387386
/**

‎dpnp/backend/kernels/dpnp_krnl_arraycreation.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -73,9 +73,9 @@ DPCTLSyclEventRef dpnp_arange_c(DPCTLSyclQueueRef q_ref,
7373

7474
event = q.submit(kernel_func);
7575

76-
event.wait();
76+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
7777

78-
return event_ref;
78+
return DPCTLEvent_Copy(event_ref);
7979
}
8080

8181
template <typename _DataType>

‎dpnp/backend/kernels/dpnp_krnl_bitwise.cpp

Lines changed: 8 additions & 5 deletions
Original file line numberDiff line numberDiff line change
@@ -68,9 +68,9 @@ DPCTLSyclEventRef dpnp_invert_c(DPCTLSyclQueueRef q_ref,
6868

6969
event = q.submit(kernel_func);
7070

71-
event.wait();
71+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
7272

73-
return event_ref;
73+
return DPCTLEvent_Copy(event_ref);
7474
}
7575

7676
template <typename _DataType>
@@ -101,6 +101,9 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
101101
fmap[DPNPFuncName::DPNP_FN_INVERT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_invert_default_c<int32_t>};
102102
fmap[DPNPFuncName::DPNP_FN_INVERT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_invert_default_c<int64_t>};
103103

104+
fmap[DPNPFuncName::DPNP_FN_INVERT_EXT][eft_INT][eft_INT] = {eft_INT, (void*)dpnp_invert_ext_c<int32_t>};
105+
fmap[DPNPFuncName::DPNP_FN_INVERT_EXT][eft_LNG][eft_LNG] = {eft_LNG, (void*)dpnp_invert_ext_c<int64_t>};
106+
104107
return;
105108
}
106109

@@ -209,7 +212,6 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
209212
cgh.parallel_for<class __name__##_strides_kernel<_DataType>>(gws, kernel_parallel_for_func); \
210213
}; \
211214
event = q.submit(kernel_func); \
212-
event.wait(); \
213215
} \
214216
else \
215217
{ \
@@ -223,9 +225,10 @@ static void func_map_init_bitwise_1arg_1type(func_map_t& fmap)
223225
cgh.parallel_for<class __name__##_kernel<_DataType>>(gws, kernel_parallel_for_func); \
224226
}; \
225227
event = q.submit(kernel_func); \
226-
event.wait(); \
227228
} \
228-
return event_ref; \
229+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
230+
\
231+
return DPCTLEvent_Copy(event_ref); \
229232
} \
230233
\
231234
template <typename _DataType> \

‎dpnp/backend/kernels/dpnp_krnl_common.cpp

Lines changed: 8 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -81,9 +81,9 @@ DPCTLSyclEventRef dpnp_astype_c(DPCTLSyclQueueRef q_ref,
8181

8282
event = q.submit(kernel_func);
8383

84-
event.wait();
84+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
8585

86-
return event_ref;
86+
return DPCTLEvent_Copy(event_ref);
8787
}
8888

8989
template <typename _DataType, typename _ResultType>
@@ -732,9 +732,9 @@ DPCTLSyclEventRef dpnp_initval_c(DPCTLSyclQueueRef q_ref,
732732

733733
sycl::event event = q.submit(kernel_func);
734734

735-
event.wait();
735+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
736736

737-
return event_ref;
737+
return DPCTLEvent_Copy(event_ref);
738738
}
739739

740740
template <typename _DataType>
@@ -756,10 +756,10 @@ void (*dpnp_initval_default_c)(void*, void*, size_t) = dpnp_initval_c<_DataType>
756756

757757
template <typename _DataType>
758758
DPCTLSyclEventRef (*dpnp_initval_ext_c)(DPCTLSyclQueueRef,
759-
void*,
760-
void*,
761-
size_t,
762-
const DPCTLEventVectorRef) = dpnp_initval_c<_DataType>;
759+
void*,
760+
void*,
761+
size_t,
762+
const DPCTLEventVectorRef) = dpnp_initval_c<_DataType>;
763763

764764
template <typename _KernelNameSpecialization>
765765
class dpnp_matmul_c_kernel;

‎dpnp/backend/kernels/dpnp_krnl_elemwise.cpp

Lines changed: 11 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -116,7 +116,6 @@
116116
gws, kernel_parallel_for_func); \
117117
}; \
118118
event = q.submit(kernel_func); \
119-
event.wait(); \
120119
} \
121120
else \
122121
{ \
@@ -142,9 +141,11 @@
142141
{ \
143142
event = q.submit(kernel_func); \
144143
} \
145-
event.wait(); \
146144
} \
147-
return event_ref; \
145+
\
146+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
147+
\
148+
return DPCTLEvent_Copy(event_ref); \
148149
} \
149150
\
150151
template <typename _DataType_input, typename _DataType_output> \
@@ -643,9 +644,9 @@ static void func_map_init_elemwise_1arg_2type(func_map_t& fmap)
643644
} \
644645
} \
645646
\
646-
event.wait(); \
647+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
647648
\
648-
return event_ref; \
649+
return DPCTLEvent_Copy(event_ref); \
649650
} \
650651
\
651652
template <typename _DataType> \
@@ -941,6 +942,8 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
941942
\
942943
input1_it->~DPNPC_id(); \
943944
input2_it->~DPNPC_id(); \
945+
\
946+
return event_ref; \
944947
} \
945948
else if (use_strides) \
946949
{ \
@@ -969,7 +972,6 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
969972
}; \
970973
\
971974
event = q.submit(kernel_func); \
972-
event.wait(); \
973975
} \
974976
else \
975977
{ \
@@ -995,9 +997,10 @@ static void func_map_init_elemwise_1arg_1type(func_map_t& fmap)
995997
}; \
996998
event = q.submit(kernel_func); \
997999
} \
998-
event.wait(); \
9991000
} \
1000-
return event_ref; \
1001+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event); \
1002+
\
1003+
return DPCTLEvent_Copy(event_ref); \
10011004
} \
10021005
\
10031006
template <typename _DataType_output, typename _DataType_input1, typename _DataType_input2> \

‎dpnp/backend/kernels/dpnp_krnl_mathematical.cpp

Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -482,7 +482,8 @@ DPCTLSyclEventRef dpnp_ediff1d_c(DPCTLSyclQueueRef q_ref,
482482
gws, kernel_parallel_for_func);
483483
};
484484
event = q.submit(kernel_func);
485-
event.wait();
485+
486+
event_ref = reinterpret_cast<DPCTLSyclEventRef>(&event);
486487

487488
return DPCTLEvent_Copy(event_ref);
488489
}

0 commit comments

Comments
 (0)