From 405a147d03a56bfb46f535079afc339539b3bdfe Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Mon, 26 Jun 2023 16:02:11 +0300 Subject: [PATCH 1/5] Add checks for reqd_work_group_size for different dimensions and fix tests --- .../kernel_features_common.h | 64 ++++++++------ .../kernel_features_reqd_work_group_size.cpp | 83 ++++++++++++++----- ...ernel_features_speculative_compilation.cpp | 8 +- ...nel_features_sub_group_size_exceptions.cpp | 14 ++-- 4 files changed, 113 insertions(+), 56 deletions(-) diff --git a/tests/optional_kernel_features/kernel_features_common.h b/tests/optional_kernel_features/kernel_features_common.h index 37ba9e509..d04c9de43 100644 --- a/tests/optional_kernel_features/kernel_features_common.h +++ b/tests/optional_kernel_features/kernel_features_common.h @@ -650,29 +650,34 @@ void run_separate_lambda(const bool is_exception_expected, single_task_action, parallel_for_action, parallel_for_wg_action); } -template +template void run_separate_lambda_nd_range(const bool is_exception_expected, const sycl::errc errc_expected, - sycl::queue &queue, + sycl::queue& queue, LambdaItemArg separate_lambda_nd_item_arg, LambdaGroupArg separate_lambda_group_arg) { - auto parallel_for_action = [&queue, separate_lambda_nd_item_arg] { + auto range = + sycl_cts::util::get_cts_object::range::get(Size, Size, Size); + auto parallel_for_action = [&queue, separate_lambda_nd_item_arg, range] { queue - .submit([&](sycl::handler &cgh) { + .submit([&](sycl::handler& cgh) { cgh.parallel_for< kernel_separate_lambda>( - sycl::nd_range{sycl::range{1}, sycl::range{1}}, + sycl::nd_range{range, range}, separate_lambda_nd_item_arg); }) .wait(); }; - auto parallel_for_wg_action = [&queue, separate_lambda_group_arg] { + auto parallel_for_wg_action = [&queue, separate_lambda_group_arg, range] { + auto groupRange = + sycl_cts::util::get_cts_object::range::get(1, 1, 1); queue - .submit([&](sycl::handler &cgh) { + .submit([&](sycl::handler& cgh) { cgh.parallel_for_work_group< kernel_separate_lambda>( - sycl::range{1}, sycl::range{1}, separate_lambda_group_arg); + groupRange, range, separate_lambda_group_arg); }) .wait(); }; @@ -805,22 +810,27 @@ void run_functor(const bool is_exception_expected, single_task_action, parallel_for_action, parallel_for_wg_action); } -template +template void run_functor_nd_range(const bool is_exception_expected, - const sycl::errc errc_expected, sycl::queue &queue) { - auto parallel_for_action = [&queue] { + const sycl::errc errc_expected, sycl::queue& queue) { + auto range = + sycl_cts::util::get_cts_object::range::get(Size, Size, Size); + + auto parallel_for_action = [&queue, &range] { queue - .submit([&](sycl::handler &cgh) { + .submit([&](sycl::handler& cgh) { cgh.parallel_for>( - sycl::nd_range{sycl::range{1}, sycl::range{1}}, Functor{}); + sycl::nd_range{range, range}, Functor{}); }) .wait(); }; - auto parallel_for_wg_action = [&queue] { + auto parallel_for_wg_action = [&queue, &range] { + auto groupRange = + sycl_cts::util::get_cts_object::range::get(1, 1, 1); queue - .submit([&](sycl::handler &cgh) { + .submit([&](sycl::handler& cgh) { cgh.parallel_for_work_group>( - sycl::range{1}, sycl::range{1}, Functor{}); + groupRange, range, Functor{}); }) .wait(); }; @@ -954,27 +964,33 @@ class kernel_submission_call; single_task_action, parallel_for_action, parallel_for_wg_action); \ } -#define RUN_SUBMISSION_CALL_ND_RANGE(IS_EXCEPTION_EXPECTED, ERRC, QUEUE, \ - ATTRIBUTE, KERNEL_NAME, ...) \ +#define RUN_SUBMISSION_CALL_ND_RANGE(SIZE, D, IS_EXCEPTION_EXPECTED, ERRC, \ + QUEUE, ATTRIBUTE, KERNEL_NAME, ...) \ \ { \ auto parallel_for_action = [&QUEUE] { \ + auto range = \ + sycl_cts::util::get_cts_object::range::get(SIZE, SIZE, SIZE); \ QUEUE \ - .submit([&](sycl::handler &cgh) { \ + .submit([&](sycl::handler& cgh) { \ cgh.parallel_for< \ kernel_submission_call>( \ - sycl::nd_range{sycl::range{1}, sycl::range{1}}, \ - [=](sycl::nd_item<1>) ATTRIBUTE { __VA_ARGS__; }); \ + sycl::nd_range{range, range}, \ + [=](sycl::nd_item) ATTRIBUTE { __VA_ARGS__; }); \ }) \ .wait(); \ }; \ auto parallel_for_wg_action = [&QUEUE] { \ + auto range = \ + sycl_cts::util::get_cts_object::range::get(SIZE, SIZE, SIZE); \ + auto groupRange = \ + sycl_cts::util::get_cts_object::range::get(1, 1, 1); \ QUEUE \ - .submit([&](sycl::handler &cgh) { \ + .submit([&](sycl::handler& cgh) { \ cgh.parallel_for_work_group< \ kernel_submission_call>( \ - sycl::range{1}, sycl::range{1}, \ - [=](sycl::group<1>) ATTRIBUTE { __VA_ARGS__; }); \ + groupRange, range, \ + [=](sycl::group) ATTRIBUTE { __VA_ARGS__; }); \ }) \ .wait(); \ }; \ diff --git a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp index 0496bcf2b..4b6431e61 100644 --- a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp +++ b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp @@ -30,48 +30,87 @@ namespace kernel_features_reqd_work_group_size { using namespace sycl_cts; using namespace kernel_features_common; -template +template class Functor { public: - [[sycl::reqd_sub_group_size(N)]] void operator()(sycl::nd_item<1>) const {} - [[sycl::reqd_sub_group_size(N)]] void operator()(sycl::group<1>) const {} + [[sycl::reqd_work_group_size(N)]] void operator()(sycl::nd_item<1>) const {} + [[sycl::reqd_work_group_size(N)]] void operator()(sycl::group<1>) const {} + + [[sycl::reqd_work_group_size(N, N)]] void operator()(sycl::nd_item<2>) const { + } + [[sycl::reqd_work_group_size(N, N)]] void operator()(sycl::group<2>) const {} + + [[sycl::reqd_work_group_size(N, N, N)]] void operator()( + sycl::nd_item<3>) const {} + [[sycl::reqd_work_group_size(N, N, N)]] void operator()( + sycl::group<3>) const {} }; -template +template class kernel_reqd_wg_size; -DISABLED_FOR_TEMPLATE_TEST_CASE_SIG(ComputeCpp, hipSYCL) -("Exceptions thrown by [[reqd_work_group_size(N)]] with unsupported size", - "[kernel_features]", ((size_t N), N), 16, 4294967295)({ - using kname = kernel_reqd_wg_size; +template +void test_size() { + INFO("N = " + std::to_string(N)); + using kname = kernel_reqd_wg_size; auto queue = util::get_cts_object::queue(); auto max_wg_size = queue.get_device().get_info(); + auto max_work_item_sizes = + queue.get_device() + .get_info>(); - bool is_exception_expected = (N > max_wg_size); + bool is_exception_expected = (std::pow(N, Dimensions) > max_wg_size); + + for (int i = 0; i < Dimensions; i++) + if (max_work_item_sizes[i] < N) is_exception_expected |= true; // Set expected error code constexpr sycl::errc expected_errc = sycl::errc::kernel_not_supported; { - const auto lambda_nd_item_arg = [](sycl::nd_item<1>) - [[sycl::reqd_work_group_size(N)]] {}; - const auto lambda_group_arg = [](sycl::group<1>) - [[sycl::reqd_work_group_size(N)]] {}; - - run_separate_lambda_nd_range(is_exception_expected, expected_errc, - queue, lambda_nd_item_arg, - lambda_group_arg); + if constexpr (Dimensions == 1) { + const auto lambda_nd_item_arg_1D = + [](sycl::nd_item<1>) [[sycl::reqd_work_group_size(N)]]{}; + const auto lambda_group_arg_1D = + [](sycl::group<1>) [[sycl::reqd_work_group_size(N)]]{}; + run_separate_lambda_nd_range( + is_exception_expected, expected_errc, queue, lambda_nd_item_arg_1D, + lambda_group_arg_1D); + } else if constexpr (Dimensions == 2) { + const auto lambda_nd_item_arg_2D = + [](sycl::nd_item<2>) [[sycl::reqd_work_group_size(N, N)]]{}; + const auto lambda_group_arg_2D = + [](sycl::group<2>) [[sycl::reqd_work_group_size(N, N)]]{}; + run_separate_lambda_nd_range( + is_exception_expected, expected_errc, queue, lambda_nd_item_arg_2D, + lambda_group_arg_2D); + } else { + const auto lambda_nd_item_arg_3D = + [](sycl::nd_item<3>) [[sycl::reqd_work_group_size(N, N, N)]]{}; + const auto lambda_group_arg_3D = + [](sycl::group<3>) [[sycl::reqd_work_group_size(N, N, N)]]{}; + run_separate_lambda_nd_range( + is_exception_expected, expected_errc, queue, lambda_nd_item_arg_3D, + lambda_group_arg_3D); + } } { - run_functor_nd_range>(is_exception_expected, expected_errc, - queue); + run_functor_nd_range, N, Dimensions>( + is_exception_expected, expected_errc, queue); } { - RUN_SUBMISSION_CALL_ND_RANGE(is_exception_expected, expected_errc, queue, - [[sycl::reqd_work_group_size(N)]], kname, - NO_KERNEL_BODY); + RUN_SUBMISSION_CALL_ND_RANGE( + N, Dimensions, is_exception_expected, expected_errc, queue, + [[sycl::reqd_work_group_size(N)]], kname, NO_KERNEL_BODY); } +} + +DISABLED_FOR_TEMPLATE_TEST_CASE_SIG(ComputeCpp, hipSYCL) +("Exceptions thrown by [[reqd_work_group_size(N)]] with unsupported size", + "[kernel_features]", ((int Dimensions), Dimensions), 1, 2, 3)({ + test_size<4, Dimensions>(); + test_size<4294967295, Dimensions>(); }); } // namespace kernel_features_reqd_work_group_size diff --git a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp index 630755161..f78c9d567 100644 --- a/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp +++ b/tests/optional_kernel_features/kernel_features_speculative_compilation.cpp @@ -183,7 +183,7 @@ DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp) { RUN_SUBMISSION_CALL_ND_RANGE( - is_exception_expected, errc_expected, queue, + testing_wg_size[0], 1, is_exception_expected, errc_expected, queue, [[sycl::reqd_work_group_size(testing_wg_size[0])]], kernel_speculative<5>, NO_KERNEL_BODY); } @@ -209,7 +209,7 @@ DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp) { RUN_SUBMISSION_CALL_ND_RANGE( - is_exception_expected, errc_expected, queue, + testing_wg_size[1], 1, is_exception_expected, errc_expected, queue, [[sycl::reqd_work_group_size(testing_wg_size[1])]], kernel_speculative<6>, NO_KERNEL_BODY); } @@ -241,7 +241,7 @@ DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp) { RUN_SUBMISSION_CALL_ND_RANGE( - is_exception_expected, errc_expected, queue, + testing_sg_size[0], 1, is_exception_expected, errc_expected, queue, [[sycl::reqd_sub_group_size(testing_sg_size[0])]], kernel_speculative<7>, NO_KERNEL_BODY); } @@ -269,7 +269,7 @@ DISABLED_FOR_TEST_CASE(hipSYCL, ComputeCpp) { RUN_SUBMISSION_CALL_ND_RANGE( - is_exception_expected, errc_expected, queue, + testing_sg_size[1], 1, is_exception_expected, errc_expected, queue, [[sycl::reqd_sub_group_size(testing_sg_size[1])]], kernel_speculative<8>, NO_KERNEL_BODY); } diff --git a/tests/optional_kernel_features/kernel_features_sub_group_size_exceptions.cpp b/tests/optional_kernel_features/kernel_features_sub_group_size_exceptions.cpp index 83dbe410f..88b576891 100644 --- a/tests/optional_kernel_features/kernel_features_sub_group_size_exceptions.cpp +++ b/tests/optional_kernel_features/kernel_features_sub_group_size_exceptions.cpp @@ -63,19 +63,21 @@ DISABLED_FOR_TEMPLATE_TEST_CASE_SIG(hipSYCL, ComputeCpp) const auto separate_lambda_group_arg = [](sycl::group<1>) [[sycl::reqd_sub_group_size(N)]] {}; - run_separate_lambda_nd_range(is_exception_expected, errc_expected, - queue, separate_lambda_nd_item_arg, - separate_lambda_group_arg); + run_separate_lambda_nd_range(is_exception_expected, errc_expected, + queue, separate_lambda_nd_item_arg, + separate_lambda_group_arg); } { using FunctorT = functor_with_attribute; - run_functor_nd_range(is_exception_expected, errc_expected, queue); + run_functor_nd_range(is_exception_expected, errc_expected, + queue); } { - RUN_SUBMISSION_CALL_ND_RANGE(is_exception_expected, errc_expected, queue, - [[sycl::reqd_sub_group_size(N)]], kname, {}); + RUN_SUBMISSION_CALL_ND_RANGE(N, 1, is_exception_expected, errc_expected, + queue, [[sycl::reqd_sub_group_size(N)]], kname, + {}); } }); } // namespace kernel_features_sub_group_size From e45bce00f69000a70b75e43b56efd11620a697b4 Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Mon, 26 Jun 2023 16:18:24 +0300 Subject: [PATCH 2/5] Fix clang-format --- .../kernel_features_reqd_work_group_size.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp index 4b6431e61..a76fda4f0 100644 --- a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp +++ b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp @@ -71,25 +71,25 @@ void test_size() { { if constexpr (Dimensions == 1) { const auto lambda_nd_item_arg_1D = - [](sycl::nd_item<1>) [[sycl::reqd_work_group_size(N)]]{}; - const auto lambda_group_arg_1D = - [](sycl::group<1>) [[sycl::reqd_work_group_size(N)]]{}; + [](sycl::nd_item<1>) [[sycl::reqd_work_group_size(N)]] {}; + const auto lambda_group_arg_1D = [](sycl::group<1>) + [[sycl::reqd_work_group_size(N)]] {}; run_separate_lambda_nd_range( is_exception_expected, expected_errc, queue, lambda_nd_item_arg_1D, lambda_group_arg_1D); } else if constexpr (Dimensions == 2) { const auto lambda_nd_item_arg_2D = - [](sycl::nd_item<2>) [[sycl::reqd_work_group_size(N, N)]]{}; + [](sycl::nd_item<2>) [[sycl::reqd_work_group_size(N, N)]] {}; const auto lambda_group_arg_2D = - [](sycl::group<2>) [[sycl::reqd_work_group_size(N, N)]]{}; + [](sycl::group<2>) [[sycl::reqd_work_group_size(N, N)]] {}; run_separate_lambda_nd_range( is_exception_expected, expected_errc, queue, lambda_nd_item_arg_2D, lambda_group_arg_2D); } else { const auto lambda_nd_item_arg_3D = - [](sycl::nd_item<3>) [[sycl::reqd_work_group_size(N, N, N)]]{}; + [](sycl::nd_item<3>) [[sycl::reqd_work_group_size(N, N, N)]] {}; const auto lambda_group_arg_3D = - [](sycl::group<3>) [[sycl::reqd_work_group_size(N, N, N)]]{}; + [](sycl::group<3>) [[sycl::reqd_work_group_size(N, N, N)]] {}; run_separate_lambda_nd_range( is_exception_expected, expected_errc, queue, lambda_nd_item_arg_3D, lambda_group_arg_3D); From 9c7e8ad000081b0eb13877c7be016123820cace9 Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Mon, 26 Jun 2023 16:25:44 +0300 Subject: [PATCH 3/5] Fix CI for hipSYCL --- .../kernel_features_reqd_work_group_size.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp index a76fda4f0..257e3ea90 100644 --- a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp +++ b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp @@ -49,6 +49,8 @@ class Functor { template class kernel_reqd_wg_size; +// FIXME: re-enable when max_work_item_sizes is implemented in hipsycl +#if !SYCL_CTS_COMPILING_WITH_HIPSYCL template void test_size() { INFO("N = " + std::to_string(N)); @@ -105,7 +107,7 @@ void test_size() { [[sycl::reqd_work_group_size(N)]], kname, NO_KERNEL_BODY); } } - +#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL DISABLED_FOR_TEMPLATE_TEST_CASE_SIG(ComputeCpp, hipSYCL) ("Exceptions thrown by [[reqd_work_group_size(N)]] with unsupported size", "[kernel_features]", ((int Dimensions), Dimensions), 1, 2, 3)({ From 7a8dded670d2b9c7995e08938900f5eeabb498c6 Mon Sep 17 00:00:00 2001 From: "Kopylova, NataliaX" Date: Mon, 26 Jun 2023 16:28:42 +0300 Subject: [PATCH 4/5] Fix CI for ComputeCpp --- .../kernel_features_reqd_work_group_size.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp index 257e3ea90..2a7eeb435 100644 --- a/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp +++ b/tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp @@ -49,8 +49,9 @@ class Functor { template class kernel_reqd_wg_size; -// FIXME: re-enable when max_work_item_sizes is implemented in hipsycl -#if !SYCL_CTS_COMPILING_WITH_HIPSYCL +// FIXME: re-enable when max_work_item_sizes is implemented in hipsycl and +// computcpp +#if !SYCL_CTS_COMPILING_WITH_HIPSYCL && !SYCL_CTS_COMPILING_WITH_COMPUTECPP template void test_size() { INFO("N = " + std::to_string(N)); @@ -107,7 +108,9 @@ void test_size() { [[sycl::reqd_work_group_size(N)]], kname, NO_KERNEL_BODY); } } -#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL +#endif // !SYCL_CTS_COMPILING_WITH_HIPSYCL && + // !SYCL_CTS_COMPILING_WITH_COMPUTECPP + DISABLED_FOR_TEMPLATE_TEST_CASE_SIG(ComputeCpp, hipSYCL) ("Exceptions thrown by [[reqd_work_group_size(N)]] with unsupported size", "[kernel_features]", ((int Dimensions), Dimensions), 1, 2, 3)({ From 977232209f481ad78c04db9fb95cddf7756cfaee Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 28 Jun 2023 16:35:15 -0700 Subject: [PATCH 5/5] Update tests/optional_kernel_features/kernel_features_common.h Co-authored-by: Ronan Keryell --- tests/optional_kernel_features/kernel_features_common.h | 1 + 1 file changed, 1 insertion(+) diff --git a/tests/optional_kernel_features/kernel_features_common.h b/tests/optional_kernel_features/kernel_features_common.h index d04c9de43..61351b164 100644 --- a/tests/optional_kernel_features/kernel_features_common.h +++ b/tests/optional_kernel_features/kernel_features_common.h @@ -964,6 +964,7 @@ class kernel_submission_call; single_task_action, parallel_for_action, parallel_for_wg_action); \ } +/// Use a macro because we need to inject C++11 attributes #define RUN_SUBMISSION_CALL_ND_RANGE(SIZE, D, IS_EXCEPTION_EXPECTED, ERRC, \ QUEUE, ATTRIBUTE, KERNEL_NAME, ...) \ \