Merge pull request #728 from kopylovanat/cover-gap_optional-features

Add checks for reqd_work_group_size for different dimensions and fix test

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 <typename KernelName, typename LambdaItemArg, typename LambdaGroupArg>
+template <typename KernelName, size_t Size = 1, int Dimensions = 1,
+ typename LambdaItemArg, typename LambdaGroupArg>
 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<Dimensions>::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<KernelName, call_type::item_arg>>(
- sycl::nd_range{sycl::range{1}, sycl::range{1}},
+ sycl::nd_range<Dimensions>{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<Dimensions>::get(1, 1, 1);
  queue
- .submit([&](sycl::handler &cgh) {
+ .submit([&](sycl::handler& cgh) {
  cgh.parallel_for_work_group<
  kernel_separate_lambda<KernelName, call_type::group_arg>>(
- 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 <typename Functor>
+template <typename Functor, size_t Size = 1, int Dimensions = 1>
 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<Dimensions>::get(Size, Size, Size);
+
+ auto parallel_for_action = [&queue, &range] {
  queue
- .submit([&](sycl::handler &cgh) {
+ .submit([&](sycl::handler& cgh) {
  cgh.parallel_for<kernel_parallel_for<Functor>>(
- sycl::nd_range{sycl::range{1}, sycl::range{1}}, Functor{});
+ sycl::nd_range<Dimensions>{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<Dimensions>::get(1, 1, 1);
  queue
- .submit([&](sycl::handler &cgh) {
+ .submit([&](sycl::handler& cgh) {
  cgh.parallel_for_work_group<kernel_parallel_for_wg<Functor>>(
- sycl::range{1}, sycl::range{1}, Functor{});
+ groupRange, range, Functor{});
  })
  .wait();
  };
@@ -954,27 +964,34 @@ 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, ...) \
+/// 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, ...) \
  \
  { \
  auto parallel_for_action = [&QUEUE] { \
+ auto range = \
+ sycl_cts::util::get_cts_object::range<D>::get(SIZE, SIZE, SIZE); \
  QUEUE \
- .submit([&](sycl::handler &cgh) { \
+ .submit([&](sycl::handler& cgh) { \
  cgh.parallel_for< \
  kernel_submission_call<KERNEL_NAME, call_type::item_arg>>( \
- sycl::nd_range{sycl::range{1}, sycl::range{1}}, \
- [=](sycl::nd_item<1>) ATTRIBUTE { __VA_ARGS__; }); \
+ sycl::nd_range<D>{range, range},  \
+ [=](sycl::nd_item<D>) ATTRIBUTE { __VA_ARGS__; }); \
  }) \
  .wait(); \
  }; \
  auto parallel_for_wg_action = [&QUEUE] { \
+ auto range = \
+ sycl_cts::util::get_cts_object::range<D>::get(SIZE, SIZE, SIZE); \
+ auto groupRange = \
+ sycl_cts::util::get_cts_object::range<D>::get(1, 1, 1); \
  QUEUE \
- .submit([&](sycl::handler &cgh) { \
+ .submit([&](sycl::handler& cgh) { \
  cgh.parallel_for_work_group< \
  kernel_submission_call<KERNEL_NAME, call_type::group_arg>>( \
- sycl::range{1}, sycl::range{1}, \
- [=](sycl::group<1>) ATTRIBUTE { __VA_ARGS__; }); \
+ groupRange, range,  \
+ [=](sycl::group<D>) ATTRIBUTE { __VA_ARGS__; }); \
  }) \
  .wait(); \
  }; \

tests/optional_kernel_features/kernel_features_reqd_work_group_size.cpp

@@ -30,48 +30,92 @@ namespace kernel_features_reqd_work_group_size {
 using namespace sycl_cts;
 using namespace kernel_features_common;
 
-template <size_t N>
+template <size_t N, int Dimensions>
 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 <size_t N>
+template <size_t N, int Dimensions>
 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<N>;
+// 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 <size_t N, int Dimensions>
+void test_size() {
+ INFO("N = " + std::to_string(N));
+ using kname = kernel_reqd_wg_size<N, Dimensions>;
  auto queue = util::get_cts_object::queue();
  auto max_wg_size =
  queue.get_device().get_info<sycl::info::device::max_work_group_size>();
+ auto max_work_item_sizes =
+ queue.get_device()
+ .get_info<sycl::info::device::max_work_item_sizes<Dimensions>>();
 
- 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<kname>(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<kname, N, Dimensions>(
+ 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<kname, N, Dimensions>(
+ 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<kname, N, Dimensions>(
+ is_exception_expected, expected_errc, queue, lambda_nd_item_arg_3D,
+ lambda_group_arg_3D);
+ }
  }
  {
- run_functor_nd_range<Functor<N>>(is_exception_expected, expected_errc,
-   queue);
+ run_functor_nd_range<Functor<N, Dimensions>, 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);
  }
+}
+#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)({
+ test_size<4, Dimensions>();
+ test_size<4294967295, Dimensions>();
 });
 
 } // namespace kernel_features_reqd_work_group_size

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);
  }

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<kname>(is_exception_expected, errc_expected,
- queue, separate_lambda_nd_item_arg,
- separate_lambda_group_arg);
+ run_separate_lambda_nd_range<kname, N>(is_exception_expected, errc_expected,
+  queue, separate_lambda_nd_item_arg,
+  separate_lambda_group_arg);
  }
 
  {
  using FunctorT = functor_with_attribute<N>;
- run_functor_nd_range<FunctorT>(is_exception_expected, errc_expected, queue);
+ run_functor_nd_range<FunctorT, N>(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