From db8bc7cc35b3d1a28644ebb1afa85a8451977c96 Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Tue, 21 Mar 2023 11:23:56 +0200 Subject: [PATCH 01/11] Create test plan for root_group extension --- test_plans/root_group.asciidoc | 119 +++++++++++++++++++++++++++++++++ 1 file changed, 119 insertions(+) create mode 100644 test_plans/root_group.asciidoc diff --git a/test_plans/root_group.asciidoc b/test_plans/root_group.asciidoc new file mode 100644 index 000000000..4576b58f0 --- /dev/null +++ b/test_plans/root_group.asciidoc @@ -0,0 +1,119 @@ +:sectnums: +:xrefstyle: short + += Test plan for sycl_ext_oneapi_root_group + +This is a test plan for the APIs described in +https://github.com/intel/llvm/blob/sycl/sycl/doc/extensions/proposed/sycl_ext_oneapi_root_group.asciidoc[sycl_ext_oneapi_root_group]. + + +== Testing scope + +=== Device coverage + +All of the tests described below are performed only on the default device that +is selected on the CTS command line. + +=== Feature test macro + +All of the tests should use `#ifdef SYCL_EXT_ONEAPI_ROOT_GROUP` so they can be skipped +if feature is not supported. + +== Tests + +* All following tests run with `Dimensions` = 1, 2, 3 +* Get root group object by +[source,c++] +---- +auto bundle = sycl::get_kernel_bundle(q.get_context()); +auto kernel = bundle.get_kernel(); +auto maxWGs = kernel.ext_oneapi_get_info(q); +auto nd_range = sycl::nd_range{global_range, local_range}; +auto props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; +q.parallel_for(nd_range, props, [=](sycl::nd_item it) { + auto root = it.ext_oneapi_get_root_group(); +}); +---- +where `global_range.size()` = `local_range.size() * maxWGs` + +* Check that `root` is of type `ext::oneapi::experimental::root_group` + +=== The `root_group` class API + +==== Members + +Check the following: + +* `id_type` is same as `id` +* `range_type` is same as `range` +* `linear_id_type` is same as `size_t` +* `id_type` is same as `id` +* `dimensions` is `int` and is equal to Dimensions +* `fence_scope` is `memory_scope` and is equal to `memory_scope::device` + +==== get_group_id + +Check that `get_group_id()` return type is `id` and return value is equal to `id()`. + +==== get_local_id + +Check that `get_local_id()` return type is `id` and return value is equal to `nd_item::get_global_id()`. + +==== get_group_range + +Check that `get_group_range()` return type is `range` and return value is equal to `1`. + +==== get_local_range + +Check that `get_local_range()` return type is `range` and return value is equal to `nd_item::get_global_range()`. + +==== get_max_local_range + +Check that `get_max_local_range()` return type is `range` and return value is equal to `get_local_range()`. + +==== get_group_linear_id + +Check that `get_group_linear_id()` return type is `size_t` and return value equals 0. + +==== get_local_linear_id + +Check that `get_local_linear_id()` return type is `size_t` and return value is equal to: + +* `get_local_id` if `Dimensions` == 1 +* `get_local_id[1] + (get_local_id[0] * get_local_range[1])` if `Dimensions` == 2 +* `get_local_id[2] + (get_local_id[1] * get_local_range[2]) + (get_local_id[0] * get_local_range[1] * get_local_range[2])` if `Dimensions` == 3 + +==== get_group_linear_range + +Check that `get_group_linear_range()` return type is `size_t` and return value is equal to `get_group_range.size()`. + +==== get_local_linear_range + +Check that `get_local_linear_range()` return type is `size_t` and return value is equal to `get_local_range.size()`. + +==== leader + +Check that `leader()` return type is `bool` and return value is equal to `get_local_id() == 0`. + +=== get_child_group + +* Get object `g` via `get_child_group(root)` +* Check that `g` is `group` +* Check that `g` equals `nd_item::get_group()` + +* Get object `sb` via `get_child_group(g)` +* Check that `g` is `sub_group` +* Check that `g` equals `nd_item::get_sub_group()` + +=== group_barrier for root_group + +For folowing as `fence_scope` parameter: + +* `sycl::memory_scope::device` +* `sycl::memory_scope::system` +* no argument, expecting to use default `root_group::fence_scope` = `sycl::memory_scope::device` + +if `fence_scope` is supported by the hardware (from query of `info::device::atomic_fence_scope_capabilities`) +check that invocation of `group_barrier(root, fence_scope)` behaves as expected: +when each of workitems writes 1 into zero-initialized global memory array +after barrier all array values read by workitems in reverse order are equal to 1. From 65078f627dae89b458477dc1fcb3d2b708f012ab Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Tue, 28 Mar 2023 14:40:49 +0300 Subject: [PATCH 02/11] Apply comments --- test_plans/root_group.asciidoc | 13 ++++++------- 1 file changed, 6 insertions(+), 7 deletions(-) diff --git a/test_plans/root_group.asciidoc b/test_plans/root_group.asciidoc index 4576b58f0..ce10caa95 100644 --- a/test_plans/root_group.asciidoc +++ b/test_plans/root_group.asciidoc @@ -28,13 +28,17 @@ if feature is not supported. auto bundle = sycl::get_kernel_bundle(q.get_context()); auto kernel = bundle.get_kernel(); auto maxWGs = kernel.ext_oneapi_get_info(q); +if (maxWGs < 1) + SKIP("maxWGs should be at least 1"); auto nd_range = sycl::nd_range{global_range, local_range}; auto props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; q.parallel_for(nd_range, props, [=](sycl::nd_item it) { auto root = it.ext_oneapi_get_root_group(); }); ---- -where `global_range.size()` = `local_range.size() * maxWGs` +where `global_range.size()` = `local_range.size() * maxWGs`. + +`maxWGs` should be at least `1`, otherwise test for this `Dimensions` is skipped. * Check that `root` is of type `ext::oneapi::experimental::root_group` @@ -47,7 +51,6 @@ Check the following: * `id_type` is same as `id` * `range_type` is same as `range` * `linear_id_type` is same as `size_t` -* `id_type` is same as `id` * `dimensions` is `int` and is equal to Dimensions * `fence_scope` is `memory_scope` and is equal to `memory_scope::device` @@ -77,11 +80,7 @@ Check that `get_group_linear_id()` return type is `size_t` and return value equa ==== get_local_linear_id -Check that `get_local_linear_id()` return type is `size_t` and return value is equal to: - -* `get_local_id` if `Dimensions` == 1 -* `get_local_id[1] + (get_local_id[0] * get_local_range[1])` if `Dimensions` == 2 -* `get_local_id[2] + (get_local_id[1] * get_local_range[2]) + (get_local_id[0] * get_local_range[1] * get_local_range[2])` if `Dimensions` == 3 +Check that `get_local_linear_id()` return type is `size_t` and the return value is equal to `nd_item::get_global_linear_id()`. ==== get_group_linear_range From b3699b2674901125414f8596a184c7334a20236b Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Tue, 28 Mar 2023 15:17:57 +0300 Subject: [PATCH 03/11] Add REQUIRE for max_num_work_group_sync --- test_plans/root_group.asciidoc | 5 +---- 1 file changed, 1 insertion(+), 4 deletions(-) diff --git a/test_plans/root_group.asciidoc b/test_plans/root_group.asciidoc index ce10caa95..c1a8c96a9 100644 --- a/test_plans/root_group.asciidoc +++ b/test_plans/root_group.asciidoc @@ -28,8 +28,7 @@ if feature is not supported. auto bundle = sycl::get_kernel_bundle(q.get_context()); auto kernel = bundle.get_kernel(); auto maxWGs = kernel.ext_oneapi_get_info(q); -if (maxWGs < 1) - SKIP("maxWGs should be at least 1"); +REQUIRE(maxWGs >= 1); auto nd_range = sycl::nd_range{global_range, local_range}; auto props = sycl::ext::oneapi::experimental::properties{sycl::ext::oneapi::experimental::use_root_sync}; q.parallel_for(nd_range, props, [=](sycl::nd_item it) { @@ -38,8 +37,6 @@ q.parallel_for(nd_range, props, [=](sycl::nd_item ---- where `global_range.size()` = `local_range.size() * maxWGs`. -`maxWGs` should be at least `1`, otherwise test for this `Dimensions` is skipped. - * Check that `root` is of type `ext::oneapi::experimental::root_group` === The `root_group` class API From e095b9da2f9e26d34d2109fe0574399dddabd4b0 Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Wed, 2 Aug 2023 10:59:07 -0700 Subject: [PATCH 04/11] Prevent signed integer overflow in `check_scan_over_group` Check that the maximum value of `T` is large enough to store the results of the scan functions. This check prevents overflow for small types. Signed-off-by: Michael Aziz --- tests/group_functions/group_scan.h | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index f98714a09..f309110c3 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -303,6 +303,10 @@ void check_scan_over_group(sycl::queue& queue, sycl::range range, OpT op, bool ret_type_e = false; bool ret_type_i = false; + if (std::sqrt(std::numeric_limits::max()) + T(init) <= range.size()) { + return; + } + std::vector local_id(range_size, 0); sycl::nd_range executionRange(range, range); From f4c4682ce869f3cb0e385e63e64d1cc80c937d4c Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Wed, 9 Aug 2023 12:53:28 -0700 Subject: [PATCH 05/11] Remove large test case from `check_scan` Signed-off-by: Michael Aziz --- tests/group_functions/group_scan.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index f309110c3..0dce89ff2 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -238,7 +238,7 @@ struct init_joint_scan_group { size_t work_group_size = work_group_range.size(); - const size_t sizes[3] = {5, work_group_size / 2, 3 * work_group_size}; + const size_t sizes[3] = {5, work_group_size / 2}; for (size_t size : sizes) { check_scan, true, I>( queue, size, executionRange, OperatorT(), op_name); From ebf19b27c6401436dc5f21122e228c3a05e2e197 Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Wed, 9 Aug 2023 13:49:29 -0700 Subject: [PATCH 06/11] Remove large test case from `check_scan` Signed-off-by: Michael Aziz --- tests/group_functions/group_scan.h | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 0dce89ff2..0edfce544 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -303,7 +303,8 @@ void check_scan_over_group(sycl::queue& queue, sycl::range range, OpT op, bool ret_type_e = false; bool ret_type_i = false; - if (std::sqrt(std::numeric_limits::max()) + T(init) <= range.size()) { + if (((range_size * (range_size + 1) / 2) + T(init)) > + std::numeric_limits::max()) { return; } From a3009997e228047863117e10b53c4d1a1a4099ec Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Thu, 10 Aug 2023 08:24:40 -0700 Subject: [PATCH 07/11] Fixes test sizes and use `REQUIRE` macro Signed-off-by: Michael Aziz --- tests/group_functions/group_scan.h | 10 ++++------ 1 file changed, 4 insertions(+), 6 deletions(-) diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 0edfce544..23c289732 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -180,7 +180,7 @@ struct joint_scan_group { sycl::nd_range executionRange(work_group_range, work_group_range); - const size_t sizes[3] = {5, work_group_size / 2, 3 * work_group_size}; + const size_t sizes[2] = {5, work_group_size}; for (size_t size : sizes) { check_scan, false>(queue, size, executionRange, OperatorT(), op_name); @@ -238,7 +238,7 @@ struct init_joint_scan_group { size_t work_group_size = work_group_range.size(); - const size_t sizes[3] = {5, work_group_size / 2}; + const size_t sizes[2] = {5, work_group_size}; for (size_t size : sizes) { check_scan, true, I>( queue, size, executionRange, OperatorT(), op_name); @@ -303,10 +303,8 @@ void check_scan_over_group(sycl::queue& queue, sycl::range range, OpT op, bool ret_type_e = false; bool ret_type_i = false; - if (((range_size * (range_size + 1) / 2) + T(init)) > - std::numeric_limits::max()) { - return; - } + REQUIRE(((range_size * (range_size + 1) / 2) + T(init)) > + std::numeric_limits::max()); std::vector local_id(range_size, 0); From e1fd71d40a12cdcaa277ec942ae703feae43479c Mon Sep 17 00:00:00 2001 From: Michael Aziz Date: Thu, 10 Aug 2023 08:34:17 -0700 Subject: [PATCH 08/11] Fix `REQUIRE` condition and update sizes Signed-off-by: Michael Aziz --- tests/group_functions/group_scan.h | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tests/group_functions/group_scan.h b/tests/group_functions/group_scan.h index 23c289732..45e8ac903 100644 --- a/tests/group_functions/group_scan.h +++ b/tests/group_functions/group_scan.h @@ -180,7 +180,7 @@ struct joint_scan_group { sycl::nd_range executionRange(work_group_range, work_group_range); - const size_t sizes[2] = {5, work_group_size}; + const size_t sizes[2] = {5, 2}; for (size_t size : sizes) { check_scan, false>(queue, size, executionRange, OperatorT(), op_name); @@ -238,7 +238,7 @@ struct init_joint_scan_group { size_t work_group_size = work_group_range.size(); - const size_t sizes[2] = {5, work_group_size}; + const size_t sizes[2] = {5, 2}; for (size_t size : sizes) { check_scan, true, I>( queue, size, executionRange, OperatorT(), op_name); @@ -303,7 +303,7 @@ void check_scan_over_group(sycl::queue& queue, sycl::range range, OpT op, bool ret_type_e = false; bool ret_type_i = false; - REQUIRE(((range_size * (range_size + 1) / 2) + T(init)) > + REQUIRE(((range_size * (range_size + 1) / 2) + T(init)) <= std::numeric_limits::max()); std::vector local_id(range_size, 0); From c10edec25b1e4846c4afa99decd3e727f37cd242 Mon Sep 17 00:00:00 2001 From: Fraser Cormack Date: Tue, 12 Sep 2023 13:37:45 +0100 Subject: [PATCH 09/11] Fix potential segfault in kernel_info test If the device doesn't report any valid sub-group sizes, it's better to halt the test than to segfault when dereferencing std::max_element. --- tests/kernel/kernel_info.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/tests/kernel/kernel_info.cpp b/tests/kernel/kernel_info.cpp index d336d9fc9..2861c25e3 100644 --- a/tests/kernel/kernel_info.cpp +++ b/tests/kernel/kernel_info.cpp @@ -203,6 +203,10 @@ TEST_CASE("Test kernel info", "[kernel]") { "device_specific::max_sub_group_size>(dev)"); { auto sub_group_sizes = dev.get_info(); + { + INFO("The device must report at least one sub-group size"); + REQUIRE(!sub_group_sizes.empty()); + } uint32_t max = *std::max_element(sub_group_sizes.begin(), sub_group_sizes.end()); INFO( From 58f80f94cbdbfb4a0dda5477a0e92bc6266a4026 Mon Sep 17 00:00:00 2001 From: "Cai, Justin" Date: Tue, 12 Sep 2023 13:19:25 -0700 Subject: [PATCH 10/11] Pass lambda with nd_item arg to run_separate_lambda_nd_range --- ...ernel_features_speculative_compilation.cpp | 36 ++++++++++--------- 1 file changed, 20 insertions(+), 16 deletions(-) diff --git a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp index e8ea75c71..ac4d271ff 100644 --- a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp +++ b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp @@ -165,14 +165,15 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (max_wg_size >= testing_wg_size[0]) { { - const auto separate_lambda_item_arg = [](sycl::item<1>) - [[sycl::reqd_work_group_size(testing_wg_size[0])]]{}; + const auto separate_lambda_nd_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_work_group_size(testing_wg_size[0])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_work_group_size(testing_wg_size[0])]]{}; run_separate_lambda_nd_range>( - is_exception_expected, errc_expected, queue, separate_lambda_item_arg, - separate_lambda_group_arg); + is_exception_expected, errc_expected, queue, + separate_lambda_nd_item_arg, separate_lambda_group_arg); } { @@ -191,14 +192,15 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (max_wg_size >= testing_wg_size[1]) { { - const auto separate_lambda_item_arg = [](sycl::item<1>) - [[sycl::reqd_work_group_size(testing_wg_size[1])]]{}; + const auto separate_lambda_nd_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_work_group_size(testing_wg_size[1])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_work_group_size(testing_wg_size[1])]]{}; run_separate_lambda_nd_range>( - is_exception_expected, errc_expected, queue, separate_lambda_item_arg, - separate_lambda_group_arg); + is_exception_expected, errc_expected, queue, + separate_lambda_nd_item_arg, separate_lambda_group_arg); } { @@ -223,14 +225,15 @@ DISABLED_FOR_TEST_CASE(hipSYCL) std::find(sg_sizes_vec.begin(), sg_sizes_vec.end(), testing_sg_size[0]); if (find_res != sg_sizes_vec.end()) { { - const auto separate_lambda_item_arg = - [](sycl::item<1>) [[sycl::reqd_sub_group_size(testing_sg_size[0])]]{}; + const auto separate_lambda_nd_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_sub_group_size(testing_sg_size[0])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_sub_group_size(testing_sg_size[0])]]{}; run_separate_lambda_nd_range>( - is_exception_expected, errc_expected, queue, separate_lambda_item_arg, - separate_lambda_group_arg); + is_exception_expected, errc_expected, queue, + separate_lambda_nd_item_arg, separate_lambda_group_arg); } { @@ -251,14 +254,15 @@ DISABLED_FOR_TEST_CASE(hipSYCL) std::find(sg_sizes_vec.begin(), sg_sizes_vec.end(), testing_sg_size[1]); if (find_res != sg_sizes_vec.end()) { { - const auto separate_lambda_item_arg = - [](sycl::item<1>) [[sycl::reqd_sub_group_size(testing_sg_size[1])]]{}; + const auto separate_lambda_nd_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_sub_group_size(testing_sg_size[1])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_sub_group_size(testing_sg_size[1])]]{}; run_separate_lambda_nd_range>( - is_exception_expected, errc_expected, queue, separate_lambda_item_arg, - separate_lambda_group_arg); + is_exception_expected, errc_expected, queue, + separate_lambda_nd_item_arg, separate_lambda_group_arg); } { From ad15b5c8135f8d96634bff7418bf0d28de0c811a Mon Sep 17 00:00:00 2001 From: "Larsen, Steffen" Date: Wed, 13 Sep 2023 04:35:07 -0700 Subject: [PATCH 11/11] Fix invalid argument for nd-range kernels This commit fixes the use of item and id arguments in nd_range parallel_for kernels. Signed-off-by: Larsen, Steffen --- .../accessor_api_common_buffer_local.h | 6 ++---- tests/accessor_legacy/accessor_api_utility.h | 17 +++++++++++++++++ .../kernel_features_speculative_compilation.cpp | 16 ++++++++++------ 3 files changed, 29 insertions(+), 10 deletions(-) diff --git a/tests/accessor_legacy/accessor_api_common_buffer_local.h b/tests/accessor_legacy/accessor_api_common_buffer_local.h index 2d8a728ac..4ea6f51a4 100644 --- a/tests/accessor_legacy/accessor_api_common_buffer_local.h +++ b/tests/accessor_legacy/accessor_api_common_buffer_local.h @@ -349,9 +349,7 @@ class buffer_accessor_get_pointer { void operator()() const { check_get_pointer(acc_mode_tag::get()); } - void operator()(sycl_id_t) const { - operator()(); - } + void operator()(common_id) const { operator()(); } private: void check_get_pointer(acc_mode_tag::generic) const { @@ -534,7 +532,7 @@ class buffer_accessor_api_rw { m_range(rng), size(size_) {} - void operator()(sycl_id_t idx) const { + void operator()(common_id idx) const { // We do not need work-item synchronization for atomic mode because of: // - load-store consistency within single work-item // - access to the different elements from different work-items diff --git a/tests/accessor_legacy/accessor_api_utility.h b/tests/accessor_legacy/accessor_api_utility.h index c42581a9a..59dfc13df 100644 --- a/tests/accessor_legacy/accessor_api_utility.h +++ b/tests/accessor_legacy/accessor_api_utility.h @@ -763,6 +763,23 @@ inline std::unique_ptr get_error_data(size_t count) { return get_buffer_input_data(count, dims, useIndexes); } +/** + * @brief Common class for nd_item and item arguments to allow for a common + * kernel definition for different parallel_for. + */ +template +struct common_id { + common_id(sycl_id_t I) : Id{I} {} + common_id(sycl::item::value> I) : Id{I} {} + common_id(sycl::nd_item::value> I) + : Id{I.get_global_id()} {} + + operator sycl_id_t() const { return Id; } + + private: + sycl_id_t Id; +}; + } // namespace #endif // SYCL_1_2_1_TESTS_ACCESSOR_ACCESSOR_API_UTILITY_H diff --git a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp index e8ea75c71..14062e156 100644 --- a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp +++ b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp @@ -165,8 +165,9 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (max_wg_size >= testing_wg_size[0]) { { - const auto separate_lambda_item_arg = [](sycl::item<1>) - [[sycl::reqd_work_group_size(testing_wg_size[0])]]{}; + const auto separate_lambda_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_work_group_size(testing_wg_size[0])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_work_group_size(testing_wg_size[0])]]{}; @@ -191,8 +192,9 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (max_wg_size >= testing_wg_size[1]) { { - const auto separate_lambda_item_arg = [](sycl::item<1>) - [[sycl::reqd_work_group_size(testing_wg_size[1])]]{}; + const auto separate_lambda_item_arg = + [](sycl::nd_item<1>) + [[sycl::reqd_work_group_size(testing_wg_size[1])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_work_group_size(testing_wg_size[1])]]{}; @@ -224,7 +226,8 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (find_res != sg_sizes_vec.end()) { { const auto separate_lambda_item_arg = - [](sycl::item<1>) [[sycl::reqd_sub_group_size(testing_sg_size[0])]]{}; + [](sycl::nd_item<1>) + [[sycl::reqd_sub_group_size(testing_sg_size[0])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_sub_group_size(testing_sg_size[0])]]{}; @@ -252,7 +255,8 @@ DISABLED_FOR_TEST_CASE(hipSYCL) if (find_res != sg_sizes_vec.end()) { { const auto separate_lambda_item_arg = - [](sycl::item<1>) [[sycl::reqd_sub_group_size(testing_sg_size[1])]]{}; + [](sycl::nd_item<1>) + [[sycl::reqd_sub_group_size(testing_sg_size[1])]] {}; const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_sub_group_size(testing_sg_size[1])]]{};