Work item queries: use thread locals, some refactoring

Author: PeterTh

include/simsycl/detail/check.hh

@@ -27,7 +27,7 @@ struct sink {
 
 #if SIMSYCL_CHECK_MODE == SIMSYCL_CHECK_NONE
 #define SIMSYCL_CHECK_MSG(CONDITION, ...)                                                                              \
-    do { (void)(CONDITION); } while(0)
+    do { simsycl::detail::sink{CONDITION, __VA_ARGS__}; } 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, ...)                                                                              \

include/simsycl/sycl.hh

@@ -45,5 +45,5 @@
 #include "sycl/usm.hh"
 #include "sycl/vec.hh"
 
-#include "sycl/khr/sub_group_queries.hh"
+#include "sycl/khr/work_item_queries.hh"
 // IWYU pragma: end_keep

include/simsycl/sycl/khr/sub_group_queries.hh

@@ -0,0 +1,46 @@
+#include <simsycl/sycl/group.hh>
+#include <simsycl/sycl/nd_item.hh>
+#include <simsycl/sycl/sub_group.hh>
+
+namespace simsycl::sycl::khr {
+
+#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES
+
+namespace detail {
+template<int Dimensions>
+thread_local std::optional<simsycl::sycl::nd_item<Dimensions>> g_khr_wi_query_this_nd_item;
+
+template<int Dimensions>
+thread_local std::optional<simsycl::sycl::group<Dimensions>> g_khr_wi_query_this_group;
+
+inline thread_local std::optional<simsycl::sycl::sub_group> g_khr_wi_query_this_sub_group;
+
+inline void khr_wi_query_check(bool val, [[maybe_unused]] const char *query_name) {
+    SIMSYCL_CHECK_MSG(val,
+        "Work item query state '%s' is not available.\n"
+        "Make sure that the query originated from a kernel launched with a sycl::nd_range argument",
+        query_name);
+}
+
+} // namespace detail
+
+template<int Dimensions>
+simsycl::sycl::nd_item<Dimensions> this_nd_item() {
+    detail::khr_wi_query_check(detail::g_khr_wi_query_this_nd_item<Dimensions>.has_value(), "this_nd_item");
+    return detail::g_khr_wi_query_this_nd_item<Dimensions>.value();
+}
+
+template<int Dimensions>
+simsycl::sycl::group<Dimensions> this_group() {
+    detail::khr_wi_query_check(detail::g_khr_wi_query_this_group<Dimensions>.has_value(), "this_group");
+    return detail::g_khr_wi_query_this_group<Dimensions>.value();
+}
+
+inline simsycl::sycl::sub_group this_sub_group() {
+    detail::khr_wi_query_check(detail::g_khr_wi_query_this_sub_group.has_value(), "this_sub_group");
+    return detail::g_khr_wi_query_this_sub_group.value();
+}
+
+#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES
+
+} // namespace simsycl::sycl::khr

include/simsycl/sycl/khr/work_item_queries.hh

@@ -1,14 +1,14 @@
+
+#include "simsycl/schedule.hh"
+#include "simsycl/detail/utils.hh"
+#include "simsycl/sycl/device.hh"
+#include "simsycl/sycl/exception.hh"
 #include "simsycl/sycl/group.hh"
-#include "simsycl/sycl/khr/sub_group_queries.hh"
+#include "simsycl/sycl/group_functions.hh" // IWYU pragma: keep
+#include "simsycl/sycl/handler.hh"         // IWYU pragma: keep
+#include "simsycl/sycl/khr/work_item_queries.hh"
 #include "simsycl/sycl/nd_item.hh"
-
-#include <simsycl/detail/utils.hh>
-#include <simsycl/schedule.hh>
-#include <simsycl/sycl/device.hh>
-#include <simsycl/sycl/exception.hh>
-#include <simsycl/sycl/group_functions.hh>
-#include <simsycl/sycl/handler.hh>
-#include <simsycl/system.hh>
+#include "simsycl/system.hh"
 
 #include <numeric>
 #include <random>