Check divergence in group/sub-group ops at end

celerity · Jan 8, 2024 · 4f3321d · 4f3321d
1 parent 4f4187b
commit 4f3321d
Show file tree

Hide file tree

Showing 4 changed files with 27 additions and 3 deletions.
diff --git a/include/simsycl/detail/group_operation_impl.hh b/include/simsycl/detail/group_operation_impl.hh
@@ -64,6 +64,8 @@ enum class group_operation_id {
     exclusive_scan,
     joint_inclusive_scan,
     inclusive_scan,
+    // for verification
+    exit,
 };
 
 // additional data required to implement and check correct use for some group operations

diff --git a/src/simsycl/schedule.cc b/src/simsycl/schedule.cc
@@ -266,6 +266,12 @@ void cooperative_for_nd_range(const sycl::device &device, const sycl::nd_range<D
 
                     try {
                         kernel(nd_item);
+                        // Add an implicit "exit" operations to groups and sub-groups to catch potential divergence on
+                        // the last group operation
+                        perform_group_operation(
+                            nd_item.get_group(), detail::group_operation_id::exit, detail::group_operation_spec{});
+                        perform_group_operation(
+                            nd_item.get_sub_group(), detail::group_operation_id::exit, detail::group_operation_spec{});
                     } catch(...) { //
                         caught_exceptions.push_back(std::current_exception());
                     }

diff --git a/test/check_tests.cc b/test/check_tests.cc
@@ -104,4 +104,14 @@ TEMPLATE_LIST_TEST_CASE("Overlapping lifetimes between host- and command-group a
 
 SIMSYCL_STOP_IGNORING_DEPRECATIONS
 
+// find a way to test this, right now the macro is set differently when compiling the lib code
+// TEST_CASE("Divergent group execution is reported", "[check][group_op]") {
+//     sycl::queue q;
+//     REQUIRE_THROWS_WITH(q.submit([&](sycl::handler &cgh) {
+//         cgh.parallel_for(sycl::nd_range<1>{2, 2}, [](sycl::nd_item<1> it) {
+//             if(it.get_global_linear_id() == 0) { group_barrier(it.get_group()); }
+//         });
+//     }),
+//         ContainsSubstring("SimSYCL check failed: op.id == new_op.id"));
+// }
 #endif
diff --git a/test/group_op_tests.cc b/test/group_op_tests.cc
@@ -9,10 +9,16 @@
 using namespace simsycl;
 
 template<sycl::Group G>
-void check_group_op_sequence(const G &g, const std::vector<detail::group_operation_id> &expected_ids) {
+void check_group_op_sequence(const G &g, std::vector<detail::group_operation_id> expected_ids) {
     auto &group_instance = detail::get_concurrent_group(g).instance;
-    CHECK(group_instance.operations.size() == expected_ids.size());
-    for(size_t i = 0; i < expected_ids.size(); ++i) { CHECK(group_instance.operations[i].id == expected_ids[i]); }
+    // remove the potential implicit "exit" operation from the end of the sequence
+    std::vector<detail::group_operation_id> actual_sequence;
+    std::transform(group_instance.operations.begin(), group_instance.operations.end(),
+        std::back_inserter(actual_sequence), [](auto &op) { return op.id; });
+    if(actual_sequence.back() == detail::group_operation_id::exit) { actual_sequence.pop_back(); }
+
+    CHECK(actual_sequence.size() == expected_ids.size());
+    for(size_t i = 0; i < expected_ids.size(); ++i) { CHECK(actual_sequence[i] == expected_ids[i]); }
 }
 
 #define REPEAT_FOR_ALL_SCHEDULES                                                                                       \