Skip to content

Commit 52fca49

Browse files
authored
Restyle statistics kernels (#116)
1 parent 8bf5ec6 commit 52fca49

File tree

1 file changed

+28
-25
lines changed

1 file changed

+28
-25
lines changed

dpnp/backend/custom_kernels_statistics.cpp

Lines changed: 28 additions & 25 deletions
Original file line numberDiff line numberDiff line change
@@ -118,20 +118,22 @@ void custom_cov_c(void* array1_in, void* result1, size_t nrows, size_t ncols)
118118
// fill lower elements
119119
cl::sycl::event event;
120120
cl::sycl::range<1> gws(nrows * nrows);
121-
event = DPNP_QUEUE.submit([&](cl::sycl::handler& cgh) {
122-
cgh.parallel_for<class custom_cov_c_kernel<_DataType> >(
123-
gws,
124-
[=](cl::sycl::id<1> global_id)
125-
{
126-
const size_t idx = global_id[0];
127-
const size_t row_idx = idx / nrows;
128-
const size_t col_idx = idx - row_idx * nrows;
129-
if (col_idx < row_idx)
130-
{
131-
result[idx] = result[col_idx * nrows + row_idx];
132-
}
133-
}); // parallel_for
134-
}); // queue.submit
121+
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
122+
const size_t idx = global_id[0];
123+
const size_t row_idx = idx / nrows;
124+
const size_t col_idx = idx - row_idx * nrows;
125+
if (col_idx < row_idx)
126+
{
127+
result[idx] = result[col_idx * nrows + row_idx];
128+
}
129+
};
130+
131+
auto kernel_func = [&](cl::sycl::handler& cgh) {
132+
cgh.parallel_for<class custom_cov_c_kernel<_DataType> >(gws, kernel_parallel_for_func);
133+
};
134+
135+
event = DPNP_QUEUE.submit(kernel_func);
136+
135137
event.wait();
136138

137139
dpnp_memory_free_c(mean);
@@ -350,18 +352,19 @@ void custom_var_c(
350352
_ResultType* squared_deviations = reinterpret_cast<_ResultType*>(dpnp_memory_alloc_c(size * sizeof(_ResultType)));
351353

352354
cl::sycl::range<1> gws(size);
353-
event = DPNP_QUEUE.submit([&](cl::sycl::handler& cgh) {
354-
cgh.parallel_for<class custom_var_c_kernel<_DataType, _ResultType> >(
355-
gws,
356-
[=](cl::sycl::id<1> global_id)
355+
auto kernel_parallel_for_func = [=](cl::sycl::id<1> global_id) {
356+
size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/
357357
{
358-
size_t i = global_id[0]; /*for (size_t i = 0; i < size; ++i)*/
359-
{
360-
_ResultType deviation = static_cast<_ResultType>(array1[i]) - mean_val;
361-
squared_deviations[i] = deviation * deviation;
362-
}
363-
}); /* parallel_for */
364-
}); /* queue.submit */
358+
_ResultType deviation = static_cast<_ResultType>(array1[i]) - mean_val;
359+
squared_deviations[i] = deviation * deviation;
360+
}
361+
};
362+
363+
auto kernel_func = [&](cl::sycl::handler& cgh) {
364+
cgh.parallel_for<class custom_var_c_kernel<_DataType, _ResultType> >(gws, kernel_parallel_for_func);
365+
};
366+
367+
event = DPNP_QUEUE.submit(kernel_func);
365368

366369
event.wait();
367370

0 commit comments

Comments
 (0)