Skip to content

Commit

Permalink
[group_functions] Insert barrier after each validation
Browse files Browse the repository at this point in the history
Otherwise, the initialization performed by the following code may change
the result buffer.

Signed-off-by: Yilong Guo <yilong.guo@intel.com>
  • Loading branch information
Nuullll committed Jul 17, 2023
1 parent 2902e86 commit 39ae3bd
Showing 1 changed file with 12 additions and 0 deletions.
12 changes: 12 additions & 0 deletions tests/group_functions/group_barrier.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -154,6 +154,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (local_acc[max_id - llid] != 1)
std::get<s::test>(group_barriers_acc[0]) = false;
// make sure we check all items before moving on
sycl::group_barrier(group);

// tests for other barriers
for (int i = 1; i < group_barrier_variants; ++i) {
Expand All @@ -172,6 +174,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (local_acc[max_id - llid] != 1)
std::get<s::test>(barrier) = false;
// make sure we check all items before moving on
sycl::group_barrier(group);

[[fallthrough]];
default:
Expand All @@ -180,6 +184,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (global_acc[max_id - llid] != 1)
std::get<s::test>(barrier) = false;
// make sure we check all items before moving on
sycl::group_barrier(group);
}
}
}
Expand All @@ -205,6 +211,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (local_acc[max_id - llid] != 1)
std::get<s::test>(sub_group_barriers_acc[0]) = false;
// make sure we check all items before moving on
sycl::group_barrier(sub_group);

// tests for other barriers
for (int i = 1; i < sub_group_barrier_variants; ++i) {
Expand All @@ -225,6 +233,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (local_acc[max_id - llid] != 1)
std::get<s::test>(barrier) = false;
// make sure we check all items before moving on
sycl::group_barrier(sub_group);

[[fallthrough]];
default:
Expand All @@ -233,6 +243,8 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1,

if (global_acc[max_id - llid] != 1)
std::get<s::test>(barrier) = false;
// make sure we check all items before moving on
sycl::group_barrier(sub_group);
}
}
}
Expand Down

0 comments on commit 39ae3bd

Please sign in to comment.