@@ -217,26 +217,14 @@ class dpnp_trace_c_kernel;
217
217
template <typename _DataType, typename _ResultType>
218
218
void dpnp_trace_c (const void * array1_in, void * result1, const size_t * shape_, const size_t ndim)
219
219
{
220
- cl::sycl::event event;
221
-
222
- if ((array1_in == nullptr ) || (result1 == nullptr ))
220
+ if ((array1_in == nullptr ) || (result1 == nullptr ) || (shape_ == nullptr ) || (ndim == 0 ))
223
221
{
224
222
return ;
225
223
}
226
224
227
225
const _DataType* array_in = reinterpret_cast <const _DataType*>(array1_in);
228
226
_ResultType* result = reinterpret_cast <_ResultType*>(result1);
229
227
230
- if (shape_ == nullptr )
231
- {
232
- return ;
233
- }
234
-
235
- if (ndim == 0 )
236
- {
237
- return ;
238
- }
239
-
240
228
size_t size = 1 ;
241
229
for (size_t i = 0 ; i < ndim - 1 ; ++i)
242
230
{
@@ -252,8 +240,8 @@ void dpnp_trace_c(const void* array1_in, void* result1, const size_t* shape_, co
252
240
auto memcpy_event = DPNP_QUEUE.memcpy (shape, shape_, ndim * sizeof (size_t ));
253
241
254
242
cl::sycl::range<1 > gws (size);
255
- auto kernel_parallel_for_func = [=](cl::sycl::id< 1 > global_id ) {
256
- size_t i = global_id [0 ];
243
+ auto kernel_parallel_for_func = [=](auto index ) {
244
+ size_t i = index [0 ];
257
245
result[i] = 0 ;
258
246
for (size_t j = 0 ; j < shape[ndim - 1 ]; ++j)
259
247
{
@@ -266,8 +254,7 @@ void dpnp_trace_c(const void* array1_in, void* result1, const size_t* shape_, co
266
254
cgh.parallel_for <class dpnp_trace_c_kernel <_DataType, _ResultType>>(gws, kernel_parallel_for_func);
267
255
};
268
256
269
- event = DPNP_QUEUE.submit (kernel_func);
270
-
257
+ auto event = DPNP_QUEUE.submit (kernel_func);
271
258
event.wait ();
272
259
273
260
dpnp_memory_free_c (shape);
0 commit comments