Improve check infrastructure, test group divergence reporting

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

Hide file tree

Showing 6 changed files with 67 additions and 35 deletions.
diff --git a/CMakeLists.txt b/CMakeLists.txt
@@ -137,6 +137,7 @@ add_library(simsycl
     src/simsycl/check.cc
     src/simsycl/context.cc
     src/simsycl/device.cc
+    src/simsycl/group_operation_impl.cc
     src/simsycl/kernel.cc
     src/simsycl/schedule.cc
     src/simsycl/platform.cc

diff --git a/include/simsycl/detail/check.hh b/include/simsycl/detail/check.hh
@@ -10,9 +10,13 @@
 #define SIMSYCL_CHECK_ABORT 4
 
 namespace simsycl::detail {
-void check_log(bool condition, const char *cond_string, std::source_location location);
-void check_throw(bool condition, const char *cond_string, std::source_location location);
-void check_abort(bool condition, const char *cond_string, std::source_location location);
+void check(
+    bool condition, const char *cond_string, std::source_location location, int default_mode, const char *message, ...);
+
+struct override_check_mode {
+    override_check_mode(int mode);
+    ~override_check_mode();
+};
 
 struct sink {
     template<typename... Args>
@@ -21,21 +25,21 @@ struct sink {
 } // namespace simsycl::detail
 
 #if SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_NONE
-#define SIMSYCL_CHECK(CONDITION)                                                                                       \
+#define SIMSYCL_CHECK_MSG(CONDITION, ...)                                                                              \
     do { (void)(CONDITION); } while(0)
-#elif SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_LOG
-#define SIMSYCL_CHECK(CONDITION)                                                                                       \
-    do { simsycl::detail::check_log(CONDITION, #CONDITION, std::source_location::current()); } while(0)
-#elif SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_THROW
-#define SIMSYCL_CHECK(CONDITION)                                                                                       \
-    do { simsycl::detail::check_throw(CONDITION, #CONDITION, std::source_location::current()); } while(0)
-#elif SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_ABORT
-#define SIMSYCL_CHECK(CONDITION)                                                                                       \
-    do { simsycl::detail::check_abort(CONDITION, #CONDITION, std::source_location::current()); } while(0)
+#elif SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_LOG || SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_THROW                             \
+    || SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_ABORT
+#define SIMSYCL_CHECK_MSG(CONDITION, ...)                                                                              \
+    do {                                                                                                               \
+        simsycl::detail::check(                                                                                        \
+            CONDITION, #CONDITION, std::source_location::current(), SIMSYCL_CHECK_MODE, __VA_ARGS__);                  \
+    } while(0)
 #else
 #error "SIMSYCL_CHECK_MODE must be SIMSYCL_CHECK_NONE, SIMSYCL_CHECK_LOG, SIMSYCL_CHECK_THROW, or SIMSYCL_CHECK_ABORT"
 #endif
 
+#define SIMSYCL_CHECK(CONDITION) SIMSYCL_CHECK_MSG(CONDITION, "")
+
 #define SIMSYCL_NOT_IMPLEMENTED                                                                                        \
     printf("SIMSYCL: Not implemented (%s:%d)\n", __FILE__, __LINE__);                                                  \
     abort();

diff --git a/include/simsycl/detail/group_operation_impl.hh b/include/simsycl/detail/group_operation_impl.hh
@@ -147,6 +147,7 @@ struct group_operation_data {
     group_operation_id id;
     size_t expected_num_work_items;
     size_t num_work_items_participating;
+    bool valid;
     std::unique_ptr<group_per_operation_data> per_op_data;
 };
 
@@ -190,6 +191,9 @@ inline detail::concurrent_sub_group &get_concurrent_group(const sycl::sub_group
     return *g.m_concurrent_group;
 }
 
+void check_group_op_validity(
+    int linear_id_in_group, const group_operation_data &new_op, group_operation_data &existing_op);
+
 // group operation function template
 
 template<typename Func>
@@ -227,10 +231,11 @@ auto perform_group_operation(G g, group_operation_id id, const Spec &spec) {
     size_t &ops_reached
         = is_sub_group_v<G> ? this_nd_item_instance.group_ops_reached : this_nd_item_instance.sub_group_ops_reached;
 
-    detail::group_operation_data new_op;
+    group_operation_data new_op;
     new_op.id = id;
     new_op.expected_num_work_items = g.get_local_range().size();
     new_op.num_work_items_participating = 1;
+    new_op.valid = true;
     new_op.per_op_data = spec.init();
 
     const size_t new_op_index = ops_reached;
@@ -243,9 +248,7 @@ auto perform_group_operation(G g, group_operation_id id, const Spec &spec) {
         SIMSYCL_CHECK(new_op_index < group_instance.operations.size() && "group operation reached in unexpected order");
 
         auto &op = group_instance.operations[ops_reached];
-        SIMSYCL_CHECK(op.id == new_op.id);
-        SIMSYCL_CHECK(op.expected_num_work_items == new_op.expected_num_work_items);
-        SIMSYCL_CHECK(op.num_work_items_participating < op.expected_num_work_items);
+        check_group_op_validity(linear_id_in_group, new_op, op);
         spec.reached(dynamic_cast<typename Spec::per_op_t &>(*op.per_op_data));
 
         op.num_work_items_participating++;
@@ -258,6 +261,7 @@ auto perform_group_operation(G g, group_operation_id id, const Spec &spec) {
         detail::yield_to_kernel_scheduler();
         // we cannot preserve a reference into `operations` across a yield since it might be resized by another item
         const auto &op = group_instance.operations[new_op_index];
+        SIMSYCL_CHECK_MSG(op.valid, "group operation invalidated by another work item");
         if(op.num_work_items_participating == op.expected_num_work_items) break;
     }
 

diff --git a/src/simsycl/check.cc b/src/simsycl/check.cc
@@ -1,7 +1,10 @@
 #include "simsycl/detail/check.hh"
 #include "simsycl/sycl/exception.hh"
 
+// TODO: use std::format/print once widely available
+#include <cassert>
 #include <iostream>
+#include <stdarg.h>
 
 namespace {
 std::string format_error(const char *cond_string, std::source_location location) {
@@ -15,18 +18,36 @@ std::string format_error(const char *cond_string, std::source_location location)
 
 namespace simsycl::detail {
 
-void check_log(bool condition, const char *cond_string, std::source_location location) {
-    if(!condition) { std::cout << format_error(cond_string, location).c_str(); }
-}
+constexpr int no_check_override = 0;
+int g_check_mode_override = no_check_override;
 
-void check_throw(bool condition, const char *cond_string, std::source_location location) {
-    if(!condition) { throw simsycl::sycl::exception(sycl::errc::invalid, format_error(cond_string, location)); }
+override_check_mode::override_check_mode(int mode) {
+    assert(g_check_mode_override == no_check_override && "check mode already overridden");
+    g_check_mode_override = mode;
 }
+override_check_mode::~override_check_mode() { g_check_mode_override = no_check_override; }
 
-void check_abort(bool condition, const char *cond_string, std::source_location location) {
+void check(bool condition, const char *cond_string, std::source_location location, int default_mode,
+    const char *message, ...) {
+    int mode = default_mode;
+    if(g_check_mode_override != no_check_override) { mode = g_check_mode_override; }
     if(!condition) {
-        std::cout << format_error(cond_string, location).c_str();
-        abort();
+        char buffer[4096];
+        va_list args;
+        va_start(args, message);
+        vsnprintf(buffer, sizeof(buffer), message, args);
+        va_end(args);
+        switch(mode) {
+            case SIMSYCL_CHECK_LOG:
+                std::cout << format_error(cond_string, location).c_str() << buffer << std::endl;
+                break;
+            case SIMSYCL_CHECK_THROW:
+                throw simsycl::sycl::exception(sycl::errc::invalid, format_error(cond_string, location) + buffer);
+            case SIMSYCL_CHECK_ABORT:
+                std::cout << format_error(cond_string, location).c_str() << buffer << std::endl;
+                abort();
+            default: assert(false && "invalid check mode");
+        }
     }
 }
 

diff --git a/test/check_tests.cc b/test/check_tests.cc
@@ -104,14 +104,4 @@ 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
@@ -1,5 +1,6 @@
 #include <catch2/catch_test_macros.hpp>
 #include <catch2/generators/catch_generators.hpp>
+#include <catch2/matchers/catch_matchers_string.hpp>
 
 #include <simsycl/schedule.hh>
 #include <sycl/sycl.hpp>
@@ -642,3 +643,14 @@ TEST_CASE("Group scans behave as expected", "[group_op][exclusive_scan_over_grou
         });
     }
 }
+
+TEST_CASE("Divergent group execution is reported", "[check][group_op]") {
+    simsycl::detail::override_check_mode check_mode(SIMSYCL_CHECK_THROW);
+    REQUIRE_THROWS_WITH(sycl::queue{}.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()); }
+        });
+    }),
+        Catch::Matchers::ContainsSubstring(
+            "group recorded operation \"barrier\", but work item #1 is trying to perform \"exit\""));
+}