diff --git a/tests/optional_kernel_features/kernel_features_common.h b/tests/optional_kernel_features/kernel_features_common.h
index 37ba9e509..61351b164 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 <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();                                                            \
     };                                                                        \
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..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
@@ -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
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<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