From df1e05d2b1fab9c71ecf2d004483408910c89140 Mon Sep 17 00:00:00 2001 From: Dmitriy Sobolev Date: Tue, 23 Apr 2024 12:12:47 -0500 Subject: [PATCH 1/4] [Tests] Improve kernel_param generation for unnamed lambda Signed-off-by: Dmitriy Sobolev --- test/support/utils.h | 13 +++++++++++-- 1 file changed, 11 insertions(+), 2 deletions(-) diff --git a/test/support/utils.h b/test/support/utils.h index 402dbbf3d4..6d56d6fa8e 100644 --- a/test/support/utils.h +++ b/test/support/utils.h @@ -48,6 +48,7 @@ #if TEST_DPCPP_BACKEND_PRESENT # include "utils_sycl.h" # include "oneapi/dpl/experimental/kt/kernel_param.h" +# include "oneapi/dpl/execution" // for oneapi::dpl::execution::DefaultKernelName #endif namespace TestUtils @@ -967,11 +968,19 @@ struct __kernel_name_with_idx template constexpr auto -get_new_kernel_params(KernelParams) +get_new_kernel_params(KernelParams params) { - return oneapi::dpl::experimental::kt::kernel_param< + auto new_params = oneapi::dpl::experimental::kt::kernel_param< KernelParams::data_per_workitem, KernelParams::workgroup_size, __kernel_name_with_idx>{}; +#if TEST_EXPLICIT_KERNEL_NAMES + return new_params; +#else + if constexpr (std::is_same_v) + return params; + else + return new_params; +#endif // TEST_EXPLICIT_KERNEL_NAMES } #endif //TEST_DPCPP_BACKEND_PRESENT From cbab2263aee6aff429bf9fbb86c66722e8139945 Mon Sep 17 00:00:00 2001 From: Dmitriy Sobolev Date: Tue, 23 Apr 2024 12:29:07 -0500 Subject: [PATCH 2/4] Add maybe_unused attribute Signed-off-by: Dmitriy Sobolev --- test/support/utils.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/test/support/utils.h b/test/support/utils.h index 6d56d6fa8e..e12777f9be 100644 --- a/test/support/utils.h +++ b/test/support/utils.h @@ -968,7 +968,7 @@ struct __kernel_name_with_idx template constexpr auto -get_new_kernel_params(KernelParams params) +get_new_kernel_params([[maybe_unused]] KernelParams params) { auto new_params = oneapi::dpl::experimental::kt::kernel_param< KernelParams::data_per_workitem, KernelParams::workgroup_size, From fd4309519f4ca2d7d926dff14694084a878792b3 Mon Sep 17 00:00:00 2001 From: Dmitriy Sobolev Date: Fri, 22 Nov 2024 14:23:50 +0000 Subject: [PATCH 3/4] Minor refactoring Signed-off-by: Dmitriy Sobolev --- test/kt/esimd_radix_sort.cpp | 18 +++++++++--------- test/kt/esimd_radix_sort_by_key.cpp | 8 ++++---- .../esimd_radix_sort_by_key_out_of_place.cpp | 8 ++++---- test/kt/esimd_radix_sort_out_of_place.cpp | 18 +++++++++--------- test/kt/single_pass_scan.cpp | 14 +++++++------- test/support/utils.h | 6 +++--- 6 files changed, 36 insertions(+), 36 deletions(-) diff --git a/test/kt/esimd_radix_sort.cpp b/test/kt/esimd_radix_sort.cpp index ecc1a18d6c..87223e429a 100755 --- a/test/kt/esimd_radix_sort.cpp +++ b/test/kt/esimd_radix_sort.cpp @@ -185,13 +185,13 @@ template (q, size, TestUtils::get_new_kernel_params<0>(param)); - test_usm(q, size, TestUtils::get_new_kernel_params<1>(param)); - test_sycl_iterators(q, size, TestUtils::get_new_kernel_params<2>(param)); - test_sycl_buffer(q, size, TestUtils::get_new_kernel_params<3>(param)); + test_usm(q, size, TestUtils::adjust_kernel_name<0>(param)); + test_usm(q, size, TestUtils::adjust_kernel_name<1>(param)); + test_sycl_iterators(q, size, TestUtils::adjust_kernel_name<2>(param)); + test_sycl_buffer(q, size, TestUtils::adjust_kernel_name<3>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, TestUtils::get_new_kernel_params<4>(param)); - test_subrange_view(q, size, TestUtils::get_new_kernel_params<5>(param)); + test_all_view(q, size, TestUtils::adjust_kernel_name<4>(param)); + test_subrange_view(q, size, TestUtils::adjust_kernel_name<5>(param)); #endif // _ENABLE_RANGES_TESTING } @@ -217,11 +217,11 @@ main() for (auto size : sort_sizes) { test_general_cases( - q, size, TestUtils::get_new_kernel_params<0>(params)); + q, size, TestUtils::adjust_kernel_name<0>(params)); test_general_cases( - q, size, TestUtils::get_new_kernel_params<1>(params)); + q, size, TestUtils::adjust_kernel_name<1>(params)); } - test_small_sizes(q, TestUtils::get_new_kernel_params<3>(params)); + test_small_sizes(q, TestUtils::adjust_kernel_name<3>(params)); } catch (const ::std::exception& exc) { diff --git a/test/kt/esimd_radix_sort_by_key.cpp b/test/kt/esimd_radix_sort_by_key.cpp index 2a48c7bef5..f045ce6b5d 100755 --- a/test/kt/esimd_radix_sort_by_key.cpp +++ b/test/kt/esimd_radix_sort_by_key.cpp @@ -103,13 +103,13 @@ int main() for (auto size : sort_sizes) { test_usm( - q, size, TestUtils::get_new_kernel_params<0>(params)); + q, size, TestUtils::adjust_kernel_name<0>(params)); test_usm( - q, size, TestUtils::get_new_kernel_params<1>(params)); + q, size, TestUtils::adjust_kernel_name<1>(params)); test_sycl_buffer( - q, size, TestUtils::get_new_kernel_params<2>(params)); + q, size, TestUtils::adjust_kernel_name<2>(params)); test_sycl_buffer( - q, size, TestUtils::get_new_kernel_params<3>(params)); + q, size, TestUtils::adjust_kernel_name<3>(params)); } } catch (const ::std::exception& exc) diff --git a/test/kt/esimd_radix_sort_by_key_out_of_place.cpp b/test/kt/esimd_radix_sort_by_key_out_of_place.cpp index d282e80d21..554f807b1f 100644 --- a/test/kt/esimd_radix_sort_by_key_out_of_place.cpp +++ b/test/kt/esimd_radix_sort_by_key_out_of_place.cpp @@ -148,13 +148,13 @@ main() for (auto size : sort_sizes) { test_usm( - q, size, TestUtils::get_new_kernel_params<0>(params)); + q, size, TestUtils::adjust_kernel_name<0>(params)); test_usm( - q, size, TestUtils::get_new_kernel_params<1>(params)); + q, size, TestUtils::adjust_kernel_name<1>(params)); test_sycl_buffer( - q, size, TestUtils::get_new_kernel_params<2>(params)); + q, size, TestUtils::adjust_kernel_name<2>(params)); test_sycl_buffer( - q, size, TestUtils::get_new_kernel_params<3>(params)); + q, size, TestUtils::adjust_kernel_name<3>(params)); } } catch (const ::std::exception& exc) diff --git a/test/kt/esimd_radix_sort_out_of_place.cpp b/test/kt/esimd_radix_sort_out_of_place.cpp index 806ac2ede6..3ee4b6f5b3 100644 --- a/test/kt/esimd_radix_sort_out_of_place.cpp +++ b/test/kt/esimd_radix_sort_out_of_place.cpp @@ -210,13 +210,13 @@ template (q, size, TestUtils::get_new_kernel_params<0>(param)); - test_usm(q, size, TestUtils::get_new_kernel_params<1>(param)); - test_sycl_iterators(q, size, TestUtils::get_new_kernel_params<2>(param)); - test_sycl_buffer(q, size, TestUtils::get_new_kernel_params<3>(param)); + test_usm(q, size, TestUtils::adjust_kernel_name<0>(param)); + test_usm(q, size, TestUtils::adjust_kernel_name<1>(param)); + test_sycl_iterators(q, size, TestUtils::adjust_kernel_name<2>(param)); + test_sycl_buffer(q, size, TestUtils::adjust_kernel_name<3>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, TestUtils::get_new_kernel_params<4>(param)); - test_subrange_view(q, size, TestUtils::get_new_kernel_params<5>(param)); + test_all_view(q, size, TestUtils::adjust_kernel_name<4>(param)); + test_subrange_view(q, size, TestUtils::adjust_kernel_name<5>(param)); #endif // _ENABLE_RANGES_TESTING } @@ -242,11 +242,11 @@ main() for (auto size : sort_sizes) { test_general_cases( - q, size, TestUtils::get_new_kernel_params<0>(params)); + q, size, TestUtils::adjust_kernel_name<0>(params)); test_general_cases( - q, size, TestUtils::get_new_kernel_params<1>(params)); + q, size, TestUtils::adjust_kernel_name<1>(params)); } - test_small_sizes(q, TestUtils::get_new_kernel_params<3>(params)); + test_small_sizes(q, TestUtils::adjust_kernel_name<3>(params)); } catch (const ::std::exception& exc) { diff --git a/test/kt/single_pass_scan.cpp b/test/kt/single_pass_scan.cpp index 860db88d2b..52f1917819 100644 --- a/test/kt/single_pass_scan.cpp +++ b/test/kt/single_pass_scan.cpp @@ -175,12 +175,12 @@ template void test_general_cases(sycl::queue q, std::size_t size, BinOp bin_op, KernelParam param) { - test_usm(q, size, bin_op, TestUtils::get_new_kernel_params<0>(param)); - test_usm(q, size, bin_op, TestUtils::get_new_kernel_params<1>(param)); - test_sycl_iterators(q, size, bin_op, TestUtils::get_new_kernel_params<2>(param)); + test_usm(q, size, bin_op, TestUtils::adjust_kernel_name<0>(param)); + test_usm(q, size, bin_op, TestUtils::adjust_kernel_name<1>(param)); + test_sycl_iterators(q, size, bin_op, TestUtils::adjust_kernel_name<2>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, bin_op, TestUtils::get_new_kernel_params<3>(param)); - test_buffer(q, size, bin_op, TestUtils::get_new_kernel_params<4>(param)); + test_all_view(q, size, bin_op, TestUtils::adjust_kernel_name<3>(param)); + test_buffer(q, size, bin_op, TestUtils::adjust_kernel_name<4>(param)); #endif } @@ -188,7 +188,7 @@ template void test_all_cases(sycl::queue q, std::size_t size, KernelParam param) { - test_general_cases(q, size, std::plus{}, TestUtils::get_new_kernel_params<0>(param)); + test_general_cases(q, size, std::plus{}, TestUtils::adjust_kernel_name<0>(param)); #if _PSTL_GROUP_REDUCTION_MULT_INT64_BROKEN static constexpr bool int64_mult_broken = std::is_integral_v && (sizeof(T) == 8); #else @@ -196,7 +196,7 @@ test_all_cases(sycl::queue q, std::size_t size, KernelParam param) #endif if constexpr (!int64_mult_broken) { - test_general_cases(q, size, std::multiplies{}, TestUtils::get_new_kernel_params<1>(param)); + test_general_cases(q, size, std::multiplies{}, TestUtils::adjust_kernel_name<1>(param)); } } diff --git a/test/support/utils.h b/test/support/utils.h index e12777f9be..3be10de01c 100644 --- a/test/support/utils.h +++ b/test/support/utils.h @@ -962,17 +962,17 @@ create_new_policy_idx(Policy&& policy) #if TEST_DPCPP_BACKEND_PRESENT template -struct __kernel_name_with_idx +struct kernel_name_with_idx { }; template constexpr auto -get_new_kernel_params([[maybe_unused]] KernelParams params) +adjust_kernel_name([[maybe_unused]] KernelParams params) { auto new_params = oneapi::dpl::experimental::kt::kernel_param< KernelParams::data_per_workitem, KernelParams::workgroup_size, - __kernel_name_with_idx>{}; + kernel_name_with_idx>{}; #if TEST_EXPLICIT_KERNEL_NAMES return new_params; #else From f81d292524694c2fd4f5ee5740faaf24cd9cc378 Mon Sep 17 00:00:00 2001 From: Dmitriy Sobolev Date: Fri, 22 Nov 2024 16:49:43 +0000 Subject: [PATCH 4/4] Align logic with make_new_policy Signed-off-by: Dmitriy Sobolev --- test/kt/esimd_radix_sort.cpp | 18 +++++++++--------- test/kt/esimd_radix_sort_by_key.cpp | 8 ++++---- .../esimd_radix_sort_by_key_out_of_place.cpp | 8 ++++---- test/kt/esimd_radix_sort_out_of_place.cpp | 18 +++++++++--------- test/kt/single_pass_scan.cpp | 14 +++++++------- test/support/utils.h | 17 ++++++----------- 6 files changed, 39 insertions(+), 44 deletions(-) diff --git a/test/kt/esimd_radix_sort.cpp b/test/kt/esimd_radix_sort.cpp index 87223e429a..a3f49f711d 100755 --- a/test/kt/esimd_radix_sort.cpp +++ b/test/kt/esimd_radix_sort.cpp @@ -185,13 +185,13 @@ template (q, size, TestUtils::adjust_kernel_name<0>(param)); - test_usm(q, size, TestUtils::adjust_kernel_name<1>(param)); - test_sycl_iterators(q, size, TestUtils::adjust_kernel_name<2>(param)); - test_sycl_buffer(q, size, TestUtils::adjust_kernel_name<3>(param)); + test_usm(q, size, TestUtils::create_new_kernel_param_idx<0>(param)); + test_usm(q, size, TestUtils::create_new_kernel_param_idx<1>(param)); + test_sycl_iterators(q, size, TestUtils::create_new_kernel_param_idx<2>(param)); + test_sycl_buffer(q, size, TestUtils::create_new_kernel_param_idx<3>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, TestUtils::adjust_kernel_name<4>(param)); - test_subrange_view(q, size, TestUtils::adjust_kernel_name<5>(param)); + test_all_view(q, size, TestUtils::create_new_kernel_param_idx<4>(param)); + test_subrange_view(q, size, TestUtils::create_new_kernel_param_idx<5>(param)); #endif // _ENABLE_RANGES_TESTING } @@ -217,11 +217,11 @@ main() for (auto size : sort_sizes) { test_general_cases( - q, size, TestUtils::adjust_kernel_name<0>(params)); + q, size, TestUtils::create_new_kernel_param_idx<0>(params)); test_general_cases( - q, size, TestUtils::adjust_kernel_name<1>(params)); + q, size, TestUtils::create_new_kernel_param_idx<1>(params)); } - test_small_sizes(q, TestUtils::adjust_kernel_name<3>(params)); + test_small_sizes(q, TestUtils::create_new_kernel_param_idx<3>(params)); } catch (const ::std::exception& exc) { diff --git a/test/kt/esimd_radix_sort_by_key.cpp b/test/kt/esimd_radix_sort_by_key.cpp index f045ce6b5d..dff02e90a7 100755 --- a/test/kt/esimd_radix_sort_by_key.cpp +++ b/test/kt/esimd_radix_sort_by_key.cpp @@ -103,13 +103,13 @@ int main() for (auto size : sort_sizes) { test_usm( - q, size, TestUtils::adjust_kernel_name<0>(params)); + q, size, TestUtils::create_new_kernel_param_idx<0>(params)); test_usm( - q, size, TestUtils::adjust_kernel_name<1>(params)); + q, size, TestUtils::create_new_kernel_param_idx<1>(params)); test_sycl_buffer( - q, size, TestUtils::adjust_kernel_name<2>(params)); + q, size, TestUtils::create_new_kernel_param_idx<2>(params)); test_sycl_buffer( - q, size, TestUtils::adjust_kernel_name<3>(params)); + q, size, TestUtils::create_new_kernel_param_idx<3>(params)); } } catch (const ::std::exception& exc) diff --git a/test/kt/esimd_radix_sort_by_key_out_of_place.cpp b/test/kt/esimd_radix_sort_by_key_out_of_place.cpp index 554f807b1f..6ebe2e2fc9 100644 --- a/test/kt/esimd_radix_sort_by_key_out_of_place.cpp +++ b/test/kt/esimd_radix_sort_by_key_out_of_place.cpp @@ -148,13 +148,13 @@ main() for (auto size : sort_sizes) { test_usm( - q, size, TestUtils::adjust_kernel_name<0>(params)); + q, size, TestUtils::create_new_kernel_param_idx<0>(params)); test_usm( - q, size, TestUtils::adjust_kernel_name<1>(params)); + q, size, TestUtils::create_new_kernel_param_idx<1>(params)); test_sycl_buffer( - q, size, TestUtils::adjust_kernel_name<2>(params)); + q, size, TestUtils::create_new_kernel_param_idx<2>(params)); test_sycl_buffer( - q, size, TestUtils::adjust_kernel_name<3>(params)); + q, size, TestUtils::create_new_kernel_param_idx<3>(params)); } } catch (const ::std::exception& exc) diff --git a/test/kt/esimd_radix_sort_out_of_place.cpp b/test/kt/esimd_radix_sort_out_of_place.cpp index 3ee4b6f5b3..cd3821e42d 100644 --- a/test/kt/esimd_radix_sort_out_of_place.cpp +++ b/test/kt/esimd_radix_sort_out_of_place.cpp @@ -210,13 +210,13 @@ template (q, size, TestUtils::adjust_kernel_name<0>(param)); - test_usm(q, size, TestUtils::adjust_kernel_name<1>(param)); - test_sycl_iterators(q, size, TestUtils::adjust_kernel_name<2>(param)); - test_sycl_buffer(q, size, TestUtils::adjust_kernel_name<3>(param)); + test_usm(q, size, TestUtils::create_new_kernel_param_idx<0>(param)); + test_usm(q, size, TestUtils::create_new_kernel_param_idx<1>(param)); + test_sycl_iterators(q, size, TestUtils::create_new_kernel_param_idx<2>(param)); + test_sycl_buffer(q, size, TestUtils::create_new_kernel_param_idx<3>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, TestUtils::adjust_kernel_name<4>(param)); - test_subrange_view(q, size, TestUtils::adjust_kernel_name<5>(param)); + test_all_view(q, size, TestUtils::create_new_kernel_param_idx<4>(param)); + test_subrange_view(q, size, TestUtils::create_new_kernel_param_idx<5>(param)); #endif // _ENABLE_RANGES_TESTING } @@ -242,11 +242,11 @@ main() for (auto size : sort_sizes) { test_general_cases( - q, size, TestUtils::adjust_kernel_name<0>(params)); + q, size, TestUtils::create_new_kernel_param_idx<0>(params)); test_general_cases( - q, size, TestUtils::adjust_kernel_name<1>(params)); + q, size, TestUtils::create_new_kernel_param_idx<1>(params)); } - test_small_sizes(q, TestUtils::adjust_kernel_name<3>(params)); + test_small_sizes(q, TestUtils::create_new_kernel_param_idx<3>(params)); } catch (const ::std::exception& exc) { diff --git a/test/kt/single_pass_scan.cpp b/test/kt/single_pass_scan.cpp index 52f1917819..f9cbe72028 100644 --- a/test/kt/single_pass_scan.cpp +++ b/test/kt/single_pass_scan.cpp @@ -175,12 +175,12 @@ template void test_general_cases(sycl::queue q, std::size_t size, BinOp bin_op, KernelParam param) { - test_usm(q, size, bin_op, TestUtils::adjust_kernel_name<0>(param)); - test_usm(q, size, bin_op, TestUtils::adjust_kernel_name<1>(param)); - test_sycl_iterators(q, size, bin_op, TestUtils::adjust_kernel_name<2>(param)); + test_usm(q, size, bin_op, TestUtils::create_new_kernel_param_idx<0>(param)); + test_usm(q, size, bin_op, TestUtils::create_new_kernel_param_idx<1>(param)); + test_sycl_iterators(q, size, bin_op, TestUtils::create_new_kernel_param_idx<2>(param)); #if _ENABLE_RANGES_TESTING - test_all_view(q, size, bin_op, TestUtils::adjust_kernel_name<3>(param)); - test_buffer(q, size, bin_op, TestUtils::adjust_kernel_name<4>(param)); + test_all_view(q, size, bin_op, TestUtils::create_new_kernel_param_idx<3>(param)); + test_buffer(q, size, bin_op, TestUtils::create_new_kernel_param_idx<4>(param)); #endif } @@ -188,7 +188,7 @@ template void test_all_cases(sycl::queue q, std::size_t size, KernelParam param) { - test_general_cases(q, size, std::plus{}, TestUtils::adjust_kernel_name<0>(param)); + test_general_cases(q, size, std::plus{}, TestUtils::create_new_kernel_param_idx<0>(param)); #if _PSTL_GROUP_REDUCTION_MULT_INT64_BROKEN static constexpr bool int64_mult_broken = std::is_integral_v && (sizeof(T) == 8); #else @@ -196,7 +196,7 @@ test_all_cases(sycl::queue q, std::size_t size, KernelParam param) #endif if constexpr (!int64_mult_broken) { - test_general_cases(q, size, std::multiplies{}, TestUtils::adjust_kernel_name<1>(param)); + test_general_cases(q, size, std::multiplies{}, TestUtils::create_new_kernel_param_idx<1>(param)); } } diff --git a/test/support/utils.h b/test/support/utils.h index 3be10de01c..666d5ca364 100644 --- a/test/support/utils.h +++ b/test/support/utils.h @@ -48,7 +48,6 @@ #if TEST_DPCPP_BACKEND_PRESENT # include "utils_sycl.h" # include "oneapi/dpl/experimental/kt/kernel_param.h" -# include "oneapi/dpl/execution" // for oneapi::dpl::execution::DefaultKernelName #endif namespace TestUtils @@ -966,20 +965,16 @@ struct kernel_name_with_idx { }; -template +template constexpr auto -adjust_kernel_name([[maybe_unused]] KernelParams params) +create_new_kernel_param_idx(KernelParam) { - auto new_params = oneapi::dpl::experimental::kt::kernel_param< - KernelParams::data_per_workitem, KernelParams::workgroup_size, - kernel_name_with_idx>{}; #if TEST_EXPLICIT_KERNEL_NAMES - return new_params; + return oneapi::dpl::experimental::kt::kernel_param>{}; #else - if constexpr (std::is_same_v) - return params; - else - return new_params; + return KernelParam{}; #endif // TEST_EXPLICIT_KERNEL_NAMES } #endif //TEST_DPCPP_BACKEND_PRESENT