From 9342b37865e1f917811d39d1e2789ca84a16ed87 Mon Sep 17 00:00:00 2001 From: Peter Thoman Date: Sun, 31 Aug 2025 20:01:44 +0200 Subject: [PATCH 1/4] Implement SYCL_KHR_WORK_ITEM_QUERIES --- .github/workflows/extension_ci.yml | 2 + CMakeLists.txt | 1 + cmake/simsycl-config.cmake.in | 1 + include/simsycl/config.hh.in | 1 + include/simsycl/sycl.hh | 2 + include/simsycl/sycl/khr/sub_group_queries.hh | 44 ++++++++++++++++ src/simsycl/schedule.cc | 49 ++++++++++++++++-- test/CMakeLists.txt | 1 + test/extensions/extensions_test.cc | 9 ++++ test/extensions/work_item_queries_test.cc | 51 +++++++++++++++++++ 10 files changed, 158 insertions(+), 3 deletions(-) create mode 100644 include/simsycl/sycl/khr/sub_group_queries.hh create mode 100644 test/extensions/work_item_queries_test.cc diff --git a/.github/workflows/extension_ci.yml b/.github/workflows/extension_ci.yml index 6ba81ad..28ff29e 100644 --- a/.github/workflows/extension_ci.yml +++ b/.github/workflows/extension_ci.yml @@ -36,6 +36,7 @@ jobs: -S ${{ github.workspace }} -DCMAKE_INSTALL_PREFIX=${{ steps.strings.outputs.install-dir }} -DSIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH=OFF + -DSIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES=OFF - name: Build SimSYCL (no extensions) run: > @@ -71,6 +72,7 @@ jobs: -S ${{ github.workspace }} -DCMAKE_INSTALL_PREFIX=${{ steps.strings.outputs.install-dir }} -DSIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH=ON + -DSIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES=ON - name: Build SimSYCL (with extensions) run: > diff --git a/CMakeLists.txt b/CMakeLists.txt index 02bb3d1..c54e412 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -69,6 +69,7 @@ set(SIMSYCL_CHECK_MODE "ABORT" CACHE STRING "Runtime assertion handling NONE|LOG # Extension options option(SIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH "Enable the SYCL_KHR_QUEUE_FLUSH extension" ON) +option(SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES "Enable the SYCL_KHR_WORK_ITEM_QUERIES extension" ON) set(CONFIG_PATH "${CMAKE_CURRENT_BINARY_DIR}/include/simsycl/config.hh") configure_file( diff --git a/cmake/simsycl-config.cmake.in b/cmake/simsycl-config.cmake.in index f203f64..3ef2efc 100644 --- a/cmake/simsycl-config.cmake.in +++ b/cmake/simsycl-config.cmake.in @@ -28,5 +28,6 @@ set(SIMSYCL_CHECK_MODE "@SIMSYCL_CHECK_MODE@") set(SIMSYCL_ENABLE_ASAN "@SIMSYCL_ENABLE_ASAN@") set(SIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH "@SIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH@") +set(SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES "@SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES@") include("${CMAKE_CURRENT_LIST_DIR}/AddToTarget.cmake") diff --git a/include/simsycl/config.hh.in b/include/simsycl/config.hh.in index 7f479df..9c1cc75 100644 --- a/include/simsycl/config.hh.in +++ b/include/simsycl/config.hh.in @@ -9,6 +9,7 @@ #cmakedefine01 SIMSYCL_FEATURE_HALF_TYPE #cmakedefine01 SIMSYCL_ENABLE_SYCL_KHR_QUEUE_FLUSH +#cmakedefine01 SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES #ifndef SIMSYCL_CHECK_MODE #define SIMSYCL_CHECK_MODE SIMSYCL_CHECK_@SIMSYCL_CHECK_MODE@ diff --git a/include/simsycl/sycl.hh b/include/simsycl/sycl.hh index 16869e4..f7623c0 100644 --- a/include/simsycl/sycl.hh +++ b/include/simsycl/sycl.hh @@ -44,4 +44,6 @@ #include "sycl/type_traits.hh" #include "sycl/usm.hh" #include "sycl/vec.hh" + +#include "sycl/khr/sub_group_queries.hh" // IWYU pragma: end_keep diff --git a/include/simsycl/sycl/khr/sub_group_queries.hh b/include/simsycl/sycl/khr/sub_group_queries.hh new file mode 100644 index 0000000..be7de77 --- /dev/null +++ b/include/simsycl/sycl/khr/sub_group_queries.hh @@ -0,0 +1,44 @@ +#include +#include +#include + +namespace simsycl::sycl::khr { + +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + +namespace detail { +template +std::optional> g_khr_wi_query_this_nd_item; + +template +std::optional> g_khr_wi_query_this_group; + +inline std::optional g_khr_wi_query_this_sub_group; +} // namespace detail + +template +simsycl::sycl::nd_item this_nd_item() { + SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_nd_item, + "Work item query state 'this_nd_item' is not available.\n" + "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); + return detail::g_khr_wi_query_this_nd_item.value(); +} + +template +simsycl::sycl::group this_group() { + SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_group, + "Work item query state 'this_group' is not available.\n" + "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); + return detail::g_khr_wi_query_this_group.value(); +} + +inline simsycl::sycl::sub_group this_sub_group() { + SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_sub_group, + "Work item query state 'this_sub_group' is not available.\n" + "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); + return detail::g_khr_wi_query_this_sub_group.value(); +} + +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + +} // namespace simsycl::sycl::khr diff --git a/src/simsycl/schedule.cc b/src/simsycl/schedule.cc index ca9ff1d..bf83758 100644 --- a/src/simsycl/schedule.cc +++ b/src/simsycl/schedule.cc @@ -1,3 +1,7 @@ +#include "simsycl/sycl/group.hh" +#include "simsycl/sycl/khr/sub_group_queries.hh" +#include "simsycl/sycl/nd_item.hh" + #include #include #include @@ -181,6 +185,23 @@ void cooperative_for_nd_range(const sycl::device &device, const sycl::nd_range concurrent_sub_groups(num_concurrent_sub_groups); std::vector num_concurrent_nd_items(num_concurrent_items); +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + std::vector *> concurrent_khr_wi_query_nd_item_ptrs(num_concurrent_items, nullptr); + + auto update_global_khr_wi_query_data = [&](int cc_g_idx = -1) { + if(cc_g_idx != -1 && concurrent_khr_wi_query_nd_item_ptrs[cc_g_idx] != nullptr) { + const auto nd_item = *concurrent_khr_wi_query_nd_item_ptrs[cc_g_idx]; + sycl::khr::detail::g_khr_wi_query_this_nd_item = nd_item; + sycl::khr::detail::g_khr_wi_query_this_group = nd_item.get_group(); + sycl::khr::detail::g_khr_wi_query_this_sub_group = nd_item.get_sub_group(); + } else { + sycl::khr::detail::g_khr_wi_query_this_nd_item = std::nullopt; + sycl::khr::detail::g_khr_wi_query_this_group = std::nullopt; + sycl::khr::detail::g_khr_wi_query_this_sub_group = std::nullopt; + } + }; +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + for(auto &cgroup : concurrent_groups) { cgroup.local_memory_allocations.resize(local_memory.size()); for(size_t i = 0; i < local_memory.size(); ++i) { @@ -220,8 +241,13 @@ void cooperative_for_nd_range(const sycl::device &device, const sycl::nd_range(local_range)) + local_id; + const auto global_id + = range.get_offset() + (group_id * sycl::id(local_range)) + local_id; // if sub-group range is not divisible by local range, the last sub-group will be smaller const auto sub_group_local_linear_range = std::min(sub_group_max_local_linear_range, @@ -265,6 +292,12 @@ void cooperative_for_nd_range(const sycl::device &device, const sycl::nd_range(1024, 64), [=](sycl::nd_item<1>) { + [[maybe_unused]] const auto item = sycl::khr::this_nd_item<1>(); + [[maybe_unused]] const auto group = sycl::khr::this_group<1>(); + [[maybe_unused]] const auto sub_group = sycl::khr::this_sub_group(); + }); + }); } diff --git a/test/extensions/work_item_queries_test.cc b/test/extensions/work_item_queries_test.cc new file mode 100644 index 0000000..c1a5b81 --- /dev/null +++ b/test/extensions/work_item_queries_test.cc @@ -0,0 +1,51 @@ +#include + +#include +#include +#include + + +using namespace simsycl; + +TEMPLATE_TEST_CASE_SIG( + "work item queries are correct if supported", "[khr][work_item_queries]", ((int Dims), Dims), 1, 2, 3) { + + #if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + + sycl::range global_range; + sycl::range local_range; + for(int d = 0; d < Dims; ++d) { + const int s = d+1; + global_range[d] = s * (2 + s); + local_range[d] = 2 + s; + } + + std::vector visited(global_range.size(), false); + sycl::queue() + .submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::nd_range(global_range, local_range), [=, &visited](sycl::nd_item it) { + const auto global_linear_id = it.get_global_linear_id(); + CHECK(global_linear_id < global_range.size()); + CHECK(!visited[global_linear_id]); + visited[global_linear_id] = true; + + CHECK(sycl::khr::this_nd_item() == it); + CHECK(sycl::khr::this_group() == it.get_group()); + CHECK(sycl::khr::this_sub_group() == it.get_sub_group()); + + group_barrier(it.get_group()); + + // check again after scheduling through group_barrier + CHECK(sycl::khr::this_nd_item() == it); + CHECK(sycl::khr::this_group() == it.get_group()); + CHECK(sycl::khr::this_sub_group() == it.get_sub_group()); + }); + }) + .wait(); + + for(size_t i = 0; i < global_range.size(); ++i) { CAPTURE(i); CHECK(visited[i]); } + + #else // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + SKIP("SYCL_KHR_WORK_ITEM_QUERIES not enabled"); + #endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +} From c0d3e831b67697e552d4e8aece9c62d66db9be5e Mon Sep 17 00:00:00 2001 From: Peter Thoman Date: Mon, 1 Sep 2025 12:50:21 +0200 Subject: [PATCH 2/4] Work item queries: use thread locals, some refactoring --- include/simsycl/detail/check.hh | 2 +- include/simsycl/sycl.hh | 2 +- include/simsycl/sycl/khr/sub_group_queries.hh | 44 ------------------ include/simsycl/sycl/khr/work_item_queries.hh | 46 +++++++++++++++++++ src/simsycl/schedule.cc | 18 ++++---- 5 files changed, 57 insertions(+), 55 deletions(-) delete mode 100644 include/simsycl/sycl/khr/sub_group_queries.hh create mode 100644 include/simsycl/sycl/khr/work_item_queries.hh diff --git a/include/simsycl/detail/check.hh b/include/simsycl/detail/check.hh index e0979d7..9565e0f 100644 --- a/include/simsycl/detail/check.hh +++ b/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, ...) \ diff --git a/include/simsycl/sycl.hh b/include/simsycl/sycl.hh index f7623c0..01d3890 100644 --- a/include/simsycl/sycl.hh +++ b/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 diff --git a/include/simsycl/sycl/khr/sub_group_queries.hh b/include/simsycl/sycl/khr/sub_group_queries.hh deleted file mode 100644 index be7de77..0000000 --- a/include/simsycl/sycl/khr/sub_group_queries.hh +++ /dev/null @@ -1,44 +0,0 @@ -#include -#include -#include - -namespace simsycl::sycl::khr { - -#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES - -namespace detail { -template -std::optional> g_khr_wi_query_this_nd_item; - -template -std::optional> g_khr_wi_query_this_group; - -inline std::optional g_khr_wi_query_this_sub_group; -} // namespace detail - -template -simsycl::sycl::nd_item this_nd_item() { - SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_nd_item, - "Work item query state 'this_nd_item' is not available.\n" - "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); - return detail::g_khr_wi_query_this_nd_item.value(); -} - -template -simsycl::sycl::group this_group() { - SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_group, - "Work item query state 'this_group' is not available.\n" - "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); - return detail::g_khr_wi_query_this_group.value(); -} - -inline simsycl::sycl::sub_group this_sub_group() { - SIMSYCL_CHECK_MSG(!!detail::g_khr_wi_query_this_sub_group, - "Work item query state 'this_sub_group' is not available.\n" - "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"); - return detail::g_khr_wi_query_this_sub_group.value(); -} - -#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES - -} // namespace simsycl::sycl::khr diff --git a/include/simsycl/sycl/khr/work_item_queries.hh b/include/simsycl/sycl/khr/work_item_queries.hh new file mode 100644 index 0000000..a850bda --- /dev/null +++ b/include/simsycl/sycl/khr/work_item_queries.hh @@ -0,0 +1,46 @@ +#include +#include +#include + +namespace simsycl::sycl::khr { + +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + +namespace detail { +template +thread_local std::optional> g_khr_wi_query_this_nd_item; + +template +thread_local std::optional> g_khr_wi_query_this_group; + +inline thread_local std::optional 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 +simsycl::sycl::nd_item this_nd_item() { + detail::khr_wi_query_check(detail::g_khr_wi_query_this_nd_item.has_value(), "this_nd_item"); + return detail::g_khr_wi_query_this_nd_item.value(); +} + +template +simsycl::sycl::group this_group() { + detail::khr_wi_query_check(detail::g_khr_wi_query_this_group.has_value(), "this_group"); + return detail::g_khr_wi_query_this_group.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 diff --git a/src/simsycl/schedule.cc b/src/simsycl/schedule.cc index bf83758..d369622 100644 --- a/src/simsycl/schedule.cc +++ b/src/simsycl/schedule.cc @@ -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 -#include -#include -#include -#include -#include -#include +#include "simsycl/system.hh" #include #include From 3b5ae850169875538836b82e253f3653f62bd381 Mon Sep 17 00:00:00 2001 From: Peter Thoman Date: Mon, 1 Sep 2025 14:24:05 +0200 Subject: [PATCH 3/4] Add error reporting unit test for SYCL_KHR_WORK_ITEM_QUERIES --- test/CMakeLists.txt | 3 ++ test/extensions/work_item_queries_test.cc | 41 +++++++++++++++++++---- 2 files changed, 38 insertions(+), 6 deletions(-) diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 1eeda80..9d4ce33 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -25,6 +25,9 @@ add_executable(tests extensions/work_item_queries_test.cc ) +# use throw check mode in the main test executable so we can have unit tests for error reporting +target_compile_definitions(tests PRIVATE SIMSYCL_CHECK_MODE=SIMSYCL_CHECK_THROW) + add_sycl_to_target(TARGET tests SIMSYCL_ALL_WARNINGS) target_link_libraries(tests PRIVATE Catch2::Catch2WithMain) diff --git a/test/extensions/work_item_queries_test.cc b/test/extensions/work_item_queries_test.cc index c1a5b81..49790fb 100644 --- a/test/extensions/work_item_queries_test.cc +++ b/test/extensions/work_item_queries_test.cc @@ -1,21 +1,23 @@ +#include "simsycl/sycl/queue.hh" #include #include #include #include +#include +using Catch::Matchers::ContainsSubstring; using namespace simsycl; TEMPLATE_TEST_CASE_SIG( "work item queries are correct if supported", "[khr][work_item_queries]", ((int Dims), Dims), 1, 2, 3) { - - #if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES sycl::range global_range; sycl::range local_range; for(int d = 0; d < Dims; ++d) { - const int s = d+1; + const int s = d + 1; global_range[d] = s * (2 + s); local_range[d] = 2 + s; } @@ -43,9 +45,36 @@ TEMPLATE_TEST_CASE_SIG( }) .wait(); - for(size_t i = 0; i < global_range.size(); ++i) { CAPTURE(i); CHECK(visited[i]); } + for(size_t i = 0; i < global_range.size(); ++i) { + CAPTURE(i); + CHECK(visited[i]); + } + +#else // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + SKIP("SYCL_KHR_WORK_ITEM_QUERIES not enabled"); +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +} + +TEST_CASE("work item queries provide useful errors", "[khr][work_item_queries]") { +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + + // outside of everything + REQUIRE_THROWS_WITH(sycl::khr::this_nd_item<1>(), ContainsSubstring("state 'this_nd_item' is not available")); + REQUIRE_THROWS_WITH(sycl::khr::this_group<1>(), ContainsSubstring("state 'this_group' is not available")); + REQUIRE_THROWS_WITH(sycl::khr::this_sub_group(), ContainsSubstring("state 'this_sub_group' is not available")); + + // in a non-nd parallel for + sycl::queue{}.submit([&](sycl::handler &cgh) { + cgh.parallel_for(sycl::range{1}, [=](sycl::item<1>) { + const char *test_str + = "Make sure that the query originated from a kernel launched with a sycl::nd_range argument"; + CHECK_THROWS_WITH(sycl::khr::this_nd_item<1>(), ContainsSubstring(test_str)); + CHECK_THROWS_WITH(sycl::khr::this_group<1>(), ContainsSubstring(test_str)); + CHECK_THROWS_WITH(sycl::khr::this_sub_group(), ContainsSubstring(test_str)); + }); + }); - #else // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +#else // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES SKIP("SYCL_KHR_WORK_ITEM_QUERIES not enabled"); - #endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES } From 81e0323e471cd3e55c9579ac3e3ba9916ba20476 Mon Sep 17 00:00:00 2001 From: Peter Thoman Date: Mon, 1 Sep 2025 15:14:53 +0200 Subject: [PATCH 4/4] Actually set the wi query feature test macro (and test that) --- include/simsycl/sycl/khr/work_item_queries.hh | 10 +++++++--- test/extensions/work_item_queries_test.cc | 7 ++++++- 2 files changed, 13 insertions(+), 4 deletions(-) diff --git a/include/simsycl/sycl/khr/work_item_queries.hh b/include/simsycl/sycl/khr/work_item_queries.hh index a850bda..4fdd41d 100644 --- a/include/simsycl/sycl/khr/work_item_queries.hh +++ b/include/simsycl/sycl/khr/work_item_queries.hh @@ -1,6 +1,10 @@ -#include -#include -#include +#include "simsycl/sycl/group.hh" +#include "simsycl/sycl/nd_item.hh" +#include "simsycl/sycl/sub_group.hh" + +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +#define SYCL_KHR_WORK_ITEM_QUERIES 1 +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES namespace simsycl::sycl::khr { diff --git a/test/extensions/work_item_queries_test.cc b/test/extensions/work_item_queries_test.cc index 49790fb..44b9c27 100644 --- a/test/extensions/work_item_queries_test.cc +++ b/test/extensions/work_item_queries_test.cc @@ -1,4 +1,3 @@ -#include "simsycl/sycl/queue.hh" #include #include @@ -10,6 +9,12 @@ using Catch::Matchers::ContainsSubstring; using namespace simsycl; +TEST_CASE("work item queries set feature test macro", "[khr][work_item_queries]") { +#if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES + CHECK(SYCL_KHR_WORK_ITEM_QUERIES == 1); +#endif // SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES +} + TEMPLATE_TEST_CASE_SIG( "work item queries are correct if supported", "[khr][work_item_queries]", ((int Dims), Dims), 1, 2, 3) { #if SIMSYCL_ENABLE_SYCL_KHR_WORK_ITEM_QUERIES