From 39ae3bddd4adeb54a4c0c3bb458f93ac0b74a1a4 Mon Sep 17 00:00:00 2001 From: Yilong Guo Date: Mon, 17 Jul 2023 11:00:46 +0800 Subject: [PATCH 1/3] [group_functions] Insert barrier after each validation Otherwise, the initialization performed by the following code may change the result buffer. Signed-off-by: Yilong Guo --- tests/group_functions/group_barrier.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/tests/group_functions/group_barrier.cpp b/tests/group_functions/group_barrier.cpp index 0dd05a452..df4997270 100644 --- a/tests/group_functions/group_barrier.cpp +++ b/tests/group_functions/group_barrier.cpp @@ -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(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) { @@ -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(barrier) = false; + // make sure we check all items before moving on + sycl::group_barrier(group); [[fallthrough]]; default: @@ -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(barrier) = false; + // make sure we check all items before moving on + sycl::group_barrier(group); } } } @@ -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(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) { @@ -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(barrier) = false; + // make sure we check all items before moving on + sycl::group_barrier(sub_group); [[fallthrough]]; default: @@ -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(barrier) = false; + // make sure we check all items before moving on + sycl::group_barrier(sub_group); } } } From d8c986c53d3f7640417e2c200c7f59da78f92f78 Mon Sep 17 00:00:00 2001 From: Yilong Guo Date: Thu, 20 Jul 2023 10:06:14 +0800 Subject: [PATCH 2/3] Add check for initialization --- tests/group_functions/group_barrier.cpp | 21 +++++++++++++++------ 1 file changed, 15 insertions(+), 6 deletions(-) diff --git a/tests/group_functions/group_barrier.cpp b/tests/group_functions/group_barrier.cpp index df4997270..72925390b 100644 --- a/tests/group_functions/group_barrier.cpp +++ b/tests/group_functions/group_barrier.cpp @@ -149,12 +149,16 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, local_acc[llid] = 0; sycl::group_barrier(group); + if (local_acc[max_id - llid] != 0) + std::get(group_barriers_acc[0]) = false; + // make sure we check all items before moving on + sycl::group_barrier(group); + local_acc[llid] = 1; sycl::group_barrier(group); if (local_acc[max_id - llid] != 1) std::get(group_barriers_acc[0]) = false; - // make sure we check all items before moving on sycl::group_barrier(group); // tests for other barriers @@ -164,7 +168,10 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if (std::get(barrier)) { local_acc[llid] = 0; global_acc[llid] = 0; + sycl::group_barrier(group); + if (local_acc[max_id - llid] != 0 || global_acc[max_id - llid] != 0) + std::get(barrier) = false; sycl::group_barrier(group); switch (std::get(barrier)) { @@ -174,7 +181,6 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if (local_acc[max_id - llid] != 1) std::get(barrier) = false; - // make sure we check all items before moving on sycl::group_barrier(group); [[fallthrough]]; @@ -184,7 +190,6 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if (global_acc[max_id - llid] != 1) std::get(barrier) = false; - // make sure we check all items before moving on sycl::group_barrier(group); } } @@ -206,12 +211,15 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, local_acc[llid] = 0; sycl::group_barrier(sub_group); + if (local_acc[max_id - llid] != 0) + std::get(sub_group_barriers_acc[0]) = false; + sycl::group_barrier(sub_group); + local_acc[llid] = 1; sycl::group_barrier(sub_group); if (local_acc[max_id - llid] != 1) std::get(sub_group_barriers_acc[0]) = false; - // make sure we check all items before moving on sycl::group_barrier(sub_group); // tests for other barriers @@ -222,7 +230,10 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, std::get(barrier)) { local_acc[llid] = 0; global_acc[llid] = 0; + sycl::group_barrier(sub_group); + if (local_acc[max_id - llid] != 0 || global_acc[max_id - llid] != 0) + std::get(barrier) = false; sycl::group_barrier(sub_group); switch (std::get(barrier)) { @@ -233,7 +244,6 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if (local_acc[max_id - llid] != 1) std::get(barrier) = false; - // make sure we check all items before moving on sycl::group_barrier(sub_group); [[fallthrough]]; @@ -243,7 +253,6 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if (global_acc[max_id - llid] != 1) std::get(barrier) = false; - // make sure we check all items before moving on sycl::group_barrier(sub_group); } } From 02a76c18141df1d0c35930392a8ff71242c901e4 Mon Sep 17 00:00:00 2001 From: Yilong Guo Date: Wed, 26 Jul 2023 15:53:42 +0800 Subject: [PATCH 3/3] Don't initialize to zero --- tests/group_functions/group_barrier.cpp | 22 ++++++++++++---------- 1 file changed, 12 insertions(+), 10 deletions(-) diff --git a/tests/group_functions/group_barrier.cpp b/tests/group_functions/group_barrier.cpp index 72925390b..d9e444236 100644 --- a/tests/group_functions/group_barrier.cpp +++ b/tests/group_functions/group_barrier.cpp @@ -146,10 +146,10 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, "fence_scope) is wrong\n"); // test of default barrier - local_acc[llid] = 0; + local_acc[llid] = llid; sycl::group_barrier(group); - if (local_acc[max_id - llid] != 0) + if (local_acc[max_id - llid] != max_id - llid) std::get(group_barriers_acc[0]) = false; // make sure we check all items before moving on sycl::group_barrier(group); @@ -166,11 +166,12 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, auto& barrier = group_barriers_acc[i]; if (std::get(barrier)) { - local_acc[llid] = 0; - global_acc[llid] = 0; + local_acc[llid] = llid; + global_acc[llid] = llid; sycl::group_barrier(group); - if (local_acc[max_id - llid] != 0 || global_acc[max_id - llid] != 0) + if (local_acc[max_id - llid] != max_id - llid || + global_acc[max_id - llid] != max_id - llid) std::get(barrier) = false; sycl::group_barrier(group); @@ -208,10 +209,10 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, "memory_scope fence_scope) is wrong\n"); // test of default barrier - local_acc[llid] = 0; + local_acc[llid] = llid; sycl::group_barrier(sub_group); - if (local_acc[max_id - llid] != 0) + if (local_acc[max_id - llid] != max_id - llid) std::get(sub_group_barriers_acc[0]) = false; sycl::group_barrier(sub_group); @@ -228,11 +229,12 @@ TEMPLATE_TEST_CASE_SIG("Group barriers", "[group_func][dim]", ((int D), D), 1, if ((sub_group.get_group_linear_id() == 0) && std::get(barrier)) { - local_acc[llid] = 0; - global_acc[llid] = 0; + local_acc[llid] = llid; + global_acc[llid] = llid; sycl::group_barrier(sub_group); - if (local_acc[max_id - llid] != 0 || global_acc[max_id - llid] != 0) + if (local_acc[max_id - llid] != max_id - llid || + global_acc[max_id - llid] != max_id - llid) std::get(barrier) = false; sycl::group_barrier(sub_group);