Skip to content

Commit 95581ee

Browse files
committed
Reduce over group
1 parent 1ec0699 commit 95581ee

File tree

1 file changed

+6
-6
lines changed

1 file changed

+6
-6
lines changed

dpnp/backend/kernels/dpnp_krnl_logic.cpp

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -251,17 +251,17 @@ DPCTLSyclEventRef dpnp_any_c(DPCTLSyclQueueRef q_ref,
251251
sycl::nd_range<1> gws(gws_range, lws_range);
252252

253253
auto kernel_parallel_for_func = [=](sycl::nd_item<1> nd_it) {
254-
auto sg = nd_it.get_sub_group();
255-
const auto max_sg_size = sg.get_max_local_range()[0];
254+
auto gr = nd_it.get_group();
255+
const auto max_gr_size = gr.get_max_local_range()[0];
256256
const size_t start =
257-
vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + sg.get_group_id()[0] * max_sg_size);
258-
const size_t end = sycl::min(start + vec_sz * max_sg_size, size);
257+
vec_sz * (nd_it.get_group(0) * nd_it.get_local_range(0) + gr.get_group_id()[0] * max_gr_size);
258+
const size_t end = sycl::min(start + vec_sz * max_gr_size, size);
259259

260260
// each work-item reduces over "vec_sz" elements in the input array
261261
bool local_reduction = sycl::joint_any_of(
262-
sg, &array_in[start], &array_in[end], [&](_DataType elem) { return elem != static_cast<_DataType>(0); });
262+
gr, &array_in[start], &array_in[end], [&](_DataType elem) { return elem != static_cast<_DataType>(0); });
263263

264-
if (sg.leader() && (local_reduction == true))
264+
if (gr.leader() && (local_reduction == true))
265265
{
266266
result[0] = true;
267267
}

0 commit comments

Comments
 (0)