diff --git a/include/sycl_ext_complex.hpp b/include/sycl_ext_complex.hpp
index b2acb76..91ffde3 100644
--- a/include/sycl_ext_complex.hpp
+++ b/include/sycl_ext_complex.hpp
@@ -375,6 +375,17 @@ _SYCL_EXT_CPLX_INLINE_VISIBILITY constexpr bool isinf(const T a) {
   return sycl::isinf(a);
 #endif
 }
+
+// To ensure loop unrolling is done when processing dimensions.
+template <size_t... Inds, class F>
+void loop_impl(std::integer_sequence<size_t, Inds...>, F &&f) {
+  (f(std::integral_constant<size_t, Inds>{}), ...);
+}
+
+template <size_t count, class F> void loop(F &&f) {
+  loop_impl(std::make_index_sequence<count>{}, std::forward<F>(f));
+}
+
 } // namespace cplex::detail
 
 ////////////////////////////////////////////////////////////////////////////////
@@ -389,12 +400,16 @@ struct is_gencomplex
                              std::is_same_v<_Tp, complex<double>> ||
                                  std::is_same_v<_Tp, complex<float>> ||
                                  std::is_same_v<_Tp, complex<sycl::half>>> {};
+template <typename _Tp>
+inline constexpr bool is_gencomplex_v = is_gencomplex<_Tp>::value;
 
 template <class _Tp>
 struct is_genfloat
     : std::integral_constant<bool, std::is_same_v<_Tp, double> ||
                                        std::is_same_v<_Tp, float> ||
                                        std::is_same_v<_Tp, sycl::half>> {};
+template <typename _Tp>
+inline constexpr bool is_genfloat_v = is_genfloat<_Tp>::value;
 
 template <class _Tp>
 class complex<_Tp, typename std::enable_if<is_genfloat<_Tp>::value>::type> {
@@ -1255,6 +1270,19 @@ _SYCL_EXT_CPLX_END_NAMESPACE_STD
 // MARRAY IMPLEMENTATION
 ////////////////////////////////////////////////////////////////////////////////
 
+_SYCL_EXT_CPLX_BEGIN_NAMESPACE_STD
+
+template <typename T> struct is_mgencomplex : std::false_type {};
+
+template <typename T, std::size_t N>
+struct is_mgencomplex<sycl::marray<T, N>>
+    : std::integral_constant<bool, sycl::ext::cplx::is_gencomplex_v<T>> {};
+
+template <typename T>
+inline constexpr bool is_mgencomplex_v = is_mgencomplex<T>::value;
+
+_SYCL_EXT_CPLX_END_NAMESPACE_STD
+
 _SYCL_MARRAY_BEGIN_NAMESPACE
 
 // marray of complex class specialisation
@@ -1679,6 +1707,497 @@ _SYCL_EXT_CPLX_INLINE_VISIBILITY
   return rtn;
 }
 
+////////////////////////////////////////////////////////////////////////////////
+// GROUP ALGORITMHS
+////////////////////////////////////////////////////////////////////////////////
+
+namespace cplex::detail {
+
+/// Helper traits to check if the type is a sycl::plus
+template <typename BinaryOperation>
+struct is_plus
+    : std::integral_constant<bool,
+                             std::is_same_v<BinaryOperation, std::plus<void>>> {
+};
+template <typename BinaryOperation>
+inline constexpr bool is_plus_v = is_plus<BinaryOperation>::value;
+
+/// Helper traits to check if the type is a sycl:multiplies
+template <typename BinaryOperation>
+struct is_multiplies
+    : std::integral_constant<
+          bool, std::is_same_v<BinaryOperation, std::multiplies<void>>> {};
+template <typename BinaryOperation>
+inline constexpr bool is_multiplies_v = is_multiplies<BinaryOperation>::value;
+
+/// Wrapper trait to check if the binary operation is supported
+template <typename BinaryOperation>
+struct is_binary_op_supported
+    : std::integral_constant<bool,
+                             (detail::is_plus<BinaryOperation>::value ||
+                              detail::is_multiplies<BinaryOperation>::value)> {
+};
+template <class BinaryOperation>
+inline constexpr bool is_binary_op_supported_v =
+    is_binary_op_supported<BinaryOperation>::value;
+
+/// Helper functions to get the init for sycl::plus binary operation when the
+/// type is a gencomplex
+template <typename T, class BinaryOperation>
+std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v<T> &&
+                  detail::is_plus_v<BinaryOperation>),
+                 T>
+get_init() {
+  return T{0, 0};
+}
+/// Helper functions to get the init for sycl::multiply binary operation when
+/// the type is a gencomplex
+template <typename T, class BinaryOperation>
+std::enable_if_t<(sycl::ext::cplx::is_gencomplex_v<T> &&
+                  detail::is_multiplies<BinaryOperation>::value),
+                 T>
+get_init() {
+  return T{1, 0};
+}
+/// Helper functions to get the init for sycl::plus binary operation when the
+/// type is a mgencomplex
+template <typename T, class BinaryOperation>
+std::enable_if_t<
+    (is_mgencomplex_v<T> && detail::is_plus<BinaryOperation>::value), T>
+get_init() {
+  using Complex = typename T::value_type;
+
+  T result;
+  std::fill(result.begin(), result.end(), Complex{0, 0});
+  return result;
+}
+/// Helper functions to get the init for sycl::multiply binary operation when
+/// the type is a mgencomplex
+template <typename T, class BinaryOperation>
+std::enable_if_t<
+    (is_mgencomplex_v<T> && detail::is_multiplies<BinaryOperation>::value), T>
+get_init() {
+  using Complex = typename T::value_type;
+
+  T result;
+  std::fill(result.begin(), result.end(), Complex{1, 0});
+  return result;
+}
+
+} // namespace cplex::detail
+
+/* REDUCE_OVER_GROUP'S OVERLOADS */
+
+/// Complex specialization
+template <typename Group, typename V, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_genfloat_v<V> &&
+              is_genfloat_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+complex<T> reduce_over_group(Group g, complex<V> x, complex<T> init,
+                             BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  complex<T> result;
+
+  result.real(sycl::reduce_over_group(g, x.real(), init.real(), binary_op));
+  result.imag(sycl::reduce_over_group(g, x.imag(), init.imag(), binary_op));
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> specialization
+template <typename Group, typename V, std::size_t N, typename T, std::size_t S,
+          class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_gencomplex_v<V> &&
+              is_gencomplex_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+sycl::marray<T, N> reduce_over_group(Group g, sycl::marray<V, N> x,
+                                     sycl::marray<T, S> init,
+                                     BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  sycl::marray<T, N> result;
+
+  cplex::detail::loop<N>([&](size_t s) {
+    result[s] = reduce_over_group(g, x[s], init[s], binary_op);
+  });
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> and Complex specialization
+template <typename Group, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              (is_gencomplex_v<T> || is_mgencomplex_v<T>)&&cplex::detail::
+                  is_binary_op_supported_v<BinaryOperation>>>
+T reduce_over_group(Group g, T x, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return reduce_over_group(g, x, init, binary_op);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/* JOINT_REDUCE'S OVERLOADS */
+
+/// Marray<Complex> and Complex specialization
+template <typename Group, typename Ptr, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              sycl::detail::is_pointer<Ptr>::value &&
+              (is_gencomplex_v<sycl::detail::remove_pointer_t<Ptr>> ||
+               is_mgencomplex_v<sycl::detail::remove_pointer_t<
+                   Ptr>>)&&(is_gencomplex_v<T> || is_mgencomplex_v<T>)&&cplex::
+                  detail::is_binary_op_supported_v<BinaryOperation>>>
+T joint_reduce(Group g, Ptr first, Ptr last, T init,
+               BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto partial = cplex::detail::get_init<T, BinaryOperation>();
+
+  sycl::detail::for_each(
+      g, first, last,
+      [&](const typename sycl::detail::remove_pointer<Ptr>::type &x) {
+        partial = binary_op(partial, x);
+      });
+
+  return reduce_over_group(g, partial, init, binary_op);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> and Complex specialization
+template <typename Group, typename Ptr, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              sycl::detail::is_pointer<Ptr>::value &&
+              (is_gencomplex_v<sycl::detail::remove_pointer_t<Ptr>> ||
+               is_mgencomplex_v<sycl::detail::remove_pointer_t<Ptr>>)&&cplex::
+                  detail::is_binary_op_supported_v<BinaryOperation>>>
+typename sycl::detail::remove_pointer_t<Ptr>
+joint_reduce(Group g, Ptr first, Ptr last, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  using T = typename sycl::detail::remove_pointer_t<Ptr>;
+
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return joint_reduce(g, first, last, init, binary_op);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/* INCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */
+
+/// Complex specialization
+template <typename Group, typename V, class BinaryOperation, typename T,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_genfloat_v<V> &&
+              is_genfloat_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+complex<T> inclusive_scan_over_group(Group g, complex<V> x,
+                                     BinaryOperation binary_op,
+                                     complex<T> init) {
+#ifdef __SYCL_DEVICE_ONLY__
+  complex<T> result;
+
+  result.real(
+      sycl::inclusive_scan_over_group(g, x.real(), binary_op, init.real()));
+  result.imag(
+      sycl::inclusive_scan_over_group(g, x.imag(), binary_op, init.imag()));
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> specialization
+template <typename Group, typename V, std::size_t N, class BinaryOperation,
+          typename T, std::size_t S,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_gencomplex_v<V> &&
+              is_gencomplex_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+sycl::marray<T, N> inclusive_scan_over_group(Group g, sycl::marray<V, N> x,
+                                             BinaryOperation binary_op,
+                                             sycl::marray<T, S> init) {
+#ifdef __SYCL_DEVICE_ONLY__
+  sycl::marray<T, N> result;
+
+  cplex::detail::loop<N>([&](size_t s) {
+    result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]);
+  });
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> and Complex specialization
+template <typename Group, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              (is_gencomplex_v<T> || is_mgencomplex_v<T>)&&cplex::detail::
+                  is_binary_op_supported_v<BinaryOperation>>>
+T inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return inclusive_scan_over_group(g, x, binary_op, init);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/* JOINT_INCLUSIVE_SCAN'S OVERLOADS */
+
+/// Complex specialization
+template <typename Group, typename InPtr, typename OutPtr,
+          class BinaryOperation, typename T,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              sycl::detail::is_pointer<InPtr>::value &&
+              sycl::detail::is_pointer<OutPtr>::value &&
+              (is_gencomplex_v<sycl::detail::remove_pointer_t<InPtr>> ||
+               is_mgencomplex_v<sycl::detail::remove_pointer_t<
+                   InPtr>>)&&(is_gencomplex_v<sycl::detail::
+                                                  remove_pointer_t<OutPtr>> ||
+                              is_mgencomplex_v<sycl::detail::remove_pointer_t<
+                                  OutPtr>>)&&(is_gencomplex_v<T> ||
+                                              is_mgencomplex_v<T>)&&cplex::
+                  detail::is_binary_op_supported_v<BinaryOperation>>>
+OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
+                            BinaryOperation binary_op, T init) {
+#ifdef __SYCL_DEVICE_ONLY__
+  std::ptrdiff_t offset = g.get_local_linear_id();
+  std::ptrdiff_t stride = g.get_local_linear_range();
+  std::ptrdiff_t N = last - first;
+
+  auto roundup = [=](const std::ptrdiff_t &v,
+                     const std::ptrdiff_t &divisor) -> std::ptrdiff_t {
+    return ((v + divisor - 1) / divisor) * divisor;
+  };
+
+  typename std::remove_const_t<typename sycl::detail::remove_pointer_t<InPtr>>
+      x;
+  typename sycl::detail::remove_pointer_t<OutPtr> carry = init;
+
+  for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
+    std::ptrdiff_t i = chunk + offset;
+
+    if (i < N)
+      x = first[i];
+
+    typename sycl::detail::remove_pointer_t<OutPtr> out =
+        inclusive_scan_over_group(g, x, binary_op, carry);
+
+    if (i < N)
+      result[i] = out;
+
+    carry = sycl::group_broadcast(g, out, stride - 1);
+  }
+  return result + N;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Complex specialization
+template <
+    typename Group, typename InPtr, typename OutPtr, class BinaryOperation,
+    typename = std::enable_if_t<
+        sycl::is_group_v<std::decay_t<Group>> &&
+        sycl::detail::is_pointer<InPtr>::value &&
+        sycl::detail::is_pointer<OutPtr>::value &&
+        (is_gencomplex_v<sycl::detail::remove_pointer_t<InPtr>> ||
+         is_mgencomplex_v<sycl::detail::remove_pointer_t<
+             InPtr>>)&&(is_gencomplex_v<sycl::detail::
+                                            remove_pointer_t<OutPtr>> ||
+                        is_mgencomplex_v<
+                            sycl::detail::remove_pointer_t<OutPtr>>)&&cplex::
+            detail::is_binary_op_supported_v<BinaryOperation>>>
+OutPtr joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
+                            BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  using T = typename sycl::detail::remove_pointer_t<InPtr>;
+
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return joint_inclusive_scan(g, first, last, result, binary_op, init);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/* EXCLUSIVE_SCAN_OVER_GROUP'S OVERLOADS */
+
+/// Complex specialization
+template <typename Group, typename V, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_genfloat_v<V> &&
+              is_genfloat_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+complex<T> exclusive_scan_over_group(Group g, complex<V> x, complex<T> init,
+                                     BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  complex<T> result;
+
+  result.real(
+      sycl::exclusive_scan_over_group(g, x.real(), init.real(), binary_op));
+  result.imag(
+      sycl::exclusive_scan_over_group(g, x.imag(), init.imag(), binary_op));
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> specialization
+template <typename Group, typename V, std::size_t N, typename T, std::size_t S,
+          class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> && is_gencomplex_v<V> &&
+              is_gencomplex_v<T> &&
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+sycl::marray<T, N> exclusive_scan_over_group(Group g, sycl::marray<V, N> x,
+                                             sycl::marray<T, S> init,
+                                             BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  sycl::marray<T, N> result;
+
+  cplex::detail::loop<N>([&](size_t s) {
+    result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op);
+  });
+
+  return result;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Marray<Complex> and Complex specialization
+template <typename Group, typename T, class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              (is_gencomplex_v<T> || is_mgencomplex_v<T>)&&cplex::detail::
+                  is_binary_op_supported_v<BinaryOperation>>>
+T exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return exclusive_scan_over_group(g, x, init, binary_op);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/* JOINT_EXCLUSIVE_SCAN'S OVERLOADS */
+
+/// Complex specialization
+template <typename Group, typename InPtr, typename OutPtr, typename T,
+          class BinaryOperation,
+          typename = std::enable_if_t<
+              sycl::is_group_v<std::decay_t<Group>> &&
+              sycl::detail::is_pointer<InPtr>::value &&
+              sycl::detail::is_pointer<OutPtr>::value &&
+              (is_gencomplex_v<sycl::detail::remove_pointer_t<InPtr>> ||
+               is_mgencomplex_v<sycl::detail::remove_pointer_t<
+                   InPtr>>)&&(is_gencomplex_v<sycl::detail::
+                                                  remove_pointer_t<OutPtr>> ||
+                              is_mgencomplex_v<sycl::detail::remove_pointer_t<
+                                  OutPtr>>)&&(is_gencomplex_v<T> ||
+                                              is_mgencomplex_v<T>)&&
+              //
+              cplex::detail::is_binary_op_supported_v<BinaryOperation>>>
+OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
+                            T init, BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  std::ptrdiff_t offset = g.get_local_linear_id();
+  std::ptrdiff_t stride = g.get_local_linear_range();
+  std::ptrdiff_t N = last - first;
+
+  auto roundup = [=](const std::ptrdiff_t &v,
+                     const std::ptrdiff_t &divisor) -> std::ptrdiff_t {
+    return ((v + divisor - 1) / divisor) * divisor;
+  };
+
+  typename std::remove_const_t<typename sycl::detail::remove_pointer_t<InPtr>>
+      x;
+  typename sycl::detail::remove_pointer_t<OutPtr> carry = init;
+
+  for (std::ptrdiff_t chunk = 0; chunk < roundup(N, stride); chunk += stride) {
+    std::ptrdiff_t i = chunk + offset;
+    if (i < N)
+      x = first[i];
+
+    typename sycl::detail::remove_pointer_t<OutPtr> out =
+        exclusive_scan_over_group(g, x, carry, binary_op);
+
+    if (i < N)
+      result[i] = out;
+
+    carry = sycl::group_broadcast(g, binary_op(out, x), stride - 1);
+  }
+  return result + N;
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
+/// Complex specialization
+template <
+    typename Group, typename InPtr, typename OutPtr, class BinaryOperation,
+    typename = std::enable_if_t<
+        sycl::is_group_v<std::decay_t<Group>> &&
+        sycl::detail::is_pointer<InPtr>::value &&
+        sycl::detail::is_pointer<OutPtr>::value &&
+        (is_gencomplex_v<sycl::detail::remove_pointer_t<InPtr>> ||
+         is_mgencomplex_v<sycl::detail::remove_pointer_t<
+             InPtr>>)&&(is_gencomplex_v<sycl::detail::
+                                            remove_pointer_t<OutPtr>> ||
+                        is_mgencomplex_v<
+                            sycl::detail::remove_pointer_t<OutPtr>>)&&cplex::
+            detail::is_binary_op_supported_v<BinaryOperation>>>
+OutPtr joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
+                            BinaryOperation binary_op) {
+#ifdef __SYCL_DEVICE_ONLY__
+  using T = typename sycl::detail::remove_pointer_t<InPtr>;
+
+  auto init = cplex::detail::get_init<T, BinaryOperation>();
+
+  return joint_exclusive_scan(g, first, last, result, init, binary_op);
+#else
+  throw sycl::exception(sycl::make_error_code(sycl::errc::runtime),
+                        "Group algorithms are not supported on host.");
+#endif
+}
+
 _SYCL_EXT_CPLX_END_NAMESPACE_STD
 
 #undef _SYCL_MARRAY_BEGIN_NAMESPACE
diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt
index 86f9467..bbd0f54 100644
--- a/tests/CMakeLists.txt
+++ b/tests/CMakeLists.txt
@@ -3,7 +3,7 @@ file(GLOB test_cases CONFIGURE_DEPENDS "*.cpp")
 foreach(test_file IN LISTS test_cases)
     if(EXISTS "${test_file}")
         get_filename_component(exe_name "${test_file}" NAME_WE)
-    
+
         add_executable(${exe_name} ${test_file})
         target_include_directories(${exe_name} PUBLIC ../include/)
         target_link_libraries(${exe_name} PRIVATE
diff --git a/tests/asin_complex.cpp b/tests/asin_complex.cpp
index a956d89..c1b84ab 100644
--- a/tests/asin_complex.cpp
+++ b/tests/asin_complex.cpp
@@ -49,7 +49,7 @@ TEMPLATE_TEST_CASE("Test complex asin", "[asin]", double, float, sycl::half) {
     }
     Q.copy(d_cplx_out, &h_cplx_out, 1).wait();
 
-    check_results(h_cplx_out, std_out, /*tol_multiplier*/ 5);
+    check_results(h_cplx_out, std_out, /*tol_multiplier*/ 6);
   }
 
   // Check cplx::complex output from host
diff --git a/tests/exclusive_scan_over_group.cpp b/tests/exclusive_scan_over_group.cpp
new file mode 100644
index 0000000..469c1ef
--- /dev/null
+++ b/tests/exclusive_scan_over_group.cpp
@@ -0,0 +1,191 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_exclusive_scan_over_group(sycl::queue q, T input,
+                                    BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(N, q);
+  auto *output_without_init = sycl::malloc_shared<V>(N, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      output_with_init[lid] = sycl::ext::cplx::exclusive_scan_over_group(
+          g, in[gid], init, binary_op);
+      output_without_init[lid] =
+          sycl::ext::cplx::exclusive_scan_over_group(g, in[gid], binary_op);
+    });
+  });
+
+  q.wait();
+
+  std::array<V, N> expected;
+  std::exclusive_scan(input.begin(), input.end(), expected.begin(), init,
+                      binary_op);
+
+  std::array<V, N> result;
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_with_init[i];
+  }
+
+  check_results(result, expected);
+
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_without_init[i];
+  }
+
+  check_results(result, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex exclusive_scan_over_group", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_exclusive_scan_over_group(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> exclusive_scan_over_group",
+                       "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_exclusive_scan_over_group(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/inclusive_scan_over_group.cpp b/tests/inclusive_scan_over_group.cpp
new file mode 100644
index 0000000..e84fe4a
--- /dev/null
+++ b/tests/inclusive_scan_over_group.cpp
@@ -0,0 +1,190 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_inclusive_scan_over_group(sycl::queue q, T input,
+                                    BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(N, q);
+  auto *output_without_init = sycl::malloc_shared<V>(N, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      output_with_init[lid] = sycl::ext::cplx::inclusive_scan_over_group(
+          g, in[gid], binary_op, init);
+      output_without_init[lid] =
+          sycl::ext::cplx::inclusive_scan_over_group(g, in[gid], binary_op);
+    });
+  });
+
+  q.wait();
+
+  std::array<V, N> expected;
+  std::inclusive_scan(input.begin(), input.end(), expected.begin(), binary_op);
+
+  std::array<V, N> result;
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_with_init[i];
+  }
+
+  check_results(result, expected);
+
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_without_init[i];
+  }
+
+  check_results(result, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex inclusive_scan_over_group", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_inclusive_scan_over_group(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> inclusive_scan_over_group",
+                       "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_inclusive_scan_over_group(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/joint_exclusive_scan.cpp b/tests/joint_exclusive_scan.cpp
new file mode 100644
index 0000000..dd3584d
--- /dev/null
+++ b/tests/joint_exclusive_scan.cpp
@@ -0,0 +1,190 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_joint_exclusive_scan(sycl::queue q, T input,
+                               BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(N, q);
+  auto *output_without_init = sycl::malloc_shared<V>(N, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      sycl::ext::cplx::joint_exclusive_scan(g, in, in + N, output_with_init,
+                                            init, binary_op);
+      sycl::ext::cplx::joint_exclusive_scan(g, in, in + N, output_without_init,
+                                            binary_op);
+    });
+  });
+
+  q.wait();
+
+  std::array<V, N> expected;
+  std::exclusive_scan(input.begin(), input.end(), expected.begin(), init,
+                      binary_op);
+
+  std::array<V, N> result;
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_with_init[i];
+  }
+
+  check_results(result, expected);
+
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_without_init[i];
+  }
+
+  check_results(result, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex joint_exclusive_scan", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_joint_exclusive_scan(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> joint_exclusive_scan", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto binary_op = BinaryOperation{};
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  if (is_type_supported<T>(q)) {
+    test_joint_exclusive_scan(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/joint_inclusive_scan.cpp b/tests/joint_inclusive_scan.cpp
new file mode 100644
index 0000000..1577fc4
--- /dev/null
+++ b/tests/joint_inclusive_scan.cpp
@@ -0,0 +1,188 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_joint_inclusive_scan(sycl::queue q, T input,
+                               BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(N, q);
+  auto *output_without_init = sycl::malloc_shared<V>(N, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      sycl::ext::cplx::joint_inclusive_scan(g, in, in + N, output_with_init,
+                                            binary_op, init);
+      sycl::ext::cplx::joint_inclusive_scan(g, in, in + N, output_without_init,
+                                            binary_op);
+    });
+  });
+
+  q.wait();
+
+  std::array<V, N> expected;
+  std::inclusive_scan(input.begin(), input.end(), expected.begin(), binary_op);
+
+  std::array<V, N> result;
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_with_init[i];
+  }
+
+  check_results(result, expected);
+
+  for (std::size_t i = 0; i < N; i++) {
+    result[i] = output_without_init[i];
+  }
+
+  check_results(result, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex joint_inclusive_scan", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_joint_inclusive_scan(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> joint_inclusive_scan", "[scan]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_joint_inclusive_scan(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/joint_reduce.cpp b/tests/joint_reduce.cpp
new file mode 100644
index 0000000..9e22158
--- /dev/null
+++ b/tests/joint_reduce.cpp
@@ -0,0 +1,178 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_joint_reduce(sycl::queue q, T input, BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  auto init = sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(1, q);
+  auto *output_without_init = sycl::malloc_shared<V>(1, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      *output_with_init =
+          sycl::ext::cplx::joint_reduce(g, in, in + N, init, binary_op);
+      *output_without_init =
+          sycl::ext::cplx::joint_reduce(g, in, in + N, binary_op);
+    });
+  });
+
+  q.wait();
+
+  const auto expected =
+      std::reduce(input.begin(), input.end(), init, binary_op);
+
+  check_results(*output_with_init, expected);
+  check_results(*output_without_init, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex joint_reduce", "[reduction]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_joint_reduce(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> joint_reduce", "[reduction]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_joint_reduce(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/reduce_over_group.cpp b/tests/reduce_over_group.cpp
new file mode 100644
index 0000000..2728c1c
--- /dev/null
+++ b/tests/reduce_over_group.cpp
@@ -0,0 +1,179 @@
+#include <array>
+#include <numeric>
+
+#include "test_helper.hpp"
+
+////////////////////////////////////////////////////////////////////////////////
+// UTILITY FUNCTIONS
+////////////////////////////////////////////////////////////////////////////////
+
+template <typename T, typename BinaryOperation>
+void test_reduce_over_group(sycl::queue q, T input, BinaryOperation binary_op) {
+  using V = typename T::value_type;
+
+  constexpr size_t N = input.size();
+
+  const auto init =
+      sycl::ext::cplx::cplex::detail::get_init<V, BinaryOperation>();
+
+  auto *in = sycl::malloc_shared<V>(N, q);
+  auto *output_with_init = sycl::malloc_shared<V>(1, q);
+  auto *output_without_init = sycl::malloc_shared<V>(1, q);
+
+  for (std::size_t i = 0; i < N; i++) {
+    in[i] = input[i];
+  }
+
+  q.submit([&](sycl::handler &cgh) {
+    cgh.parallel_for(sycl::nd_range<1>(N, N), [=](sycl::nd_item<1> it) {
+      auto gid = it.get_global_id(0);
+      auto lid = it.get_local_id(0);
+      auto g = it.get_group();
+
+      *output_with_init =
+          sycl::ext::cplx::reduce_over_group(g, in[gid], init, binary_op);
+      *output_without_init =
+          sycl::ext::cplx::reduce_over_group(g, in[gid], binary_op);
+    });
+  });
+
+  q.wait();
+
+  const auto expected =
+      std::reduce(input.begin(), input.end(), init, binary_op);
+
+  check_results(*output_with_init, expected);
+  check_results(*output_without_init, expected);
+
+  sycl::free(in, q);
+  sycl::free(output_with_init, q);
+  sycl::free(output_without_init, q);
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// COMPLEX TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test complex reduce_over_group", "[reduction]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Array = typename std::array<Complex, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Complex{1, 0}, Complex{2, 0}, Complex{3, 0}, Complex{4, 0}},
+      // Random value test
+      Array{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+            Complex{3.7, 3.7}},
+      // Repeated value test
+      Array{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+      // Negative value test
+      Array{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+            Complex{0, 0}},
+      // Large value test
+      Array{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+            Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+      // Small value test
+      Array{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+            Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+      // Edge case value test
+      Array{Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+            Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_reduce_over_group(q, test_cases, binary_op);
+  }
+}
+
+////////////////////////////////////////////////////////////////////////////////
+// MARRAY<COMPLEX> TESTS
+////////////////////////////////////////////////////////////////////////////////
+
+TEMPLATE_TEST_CASE_SIG("Test marray<complex> reduce_over_group", "[reduction]",
+                       ((typename T, std::size_t N, typename BinaryOperation),
+                        T, N, BinaryOperation),
+                       (double, 4, sycl::plus<>), (float, 4, sycl::plus<>),
+                       (sycl::half, 4, sycl::plus<>)) {
+
+  using Complex = typename sycl::ext::cplx::complex<T>;
+  using Marray = typename sycl::marray<Complex, N>;
+  using Array = typename std::array<Marray, N>;
+
+  sycl::queue q;
+
+  const auto test_cases = GENERATE(
+      // Basic value test
+      Array{Marray{Complex{1, 0}, Complex{1, 0}, Complex{1, 0}, Complex{1, 0}},
+            Marray{Complex{2, 0}, Complex{2, 0}, Complex{2, 0}, Complex{2, 0}},
+            Marray{Complex{3, 0}, Complex{3, 0}, Complex{3, 0}, Complex{3, 0}},
+            Marray{Complex{4, 0}, Complex{4, 0}, Complex{4, 0}, Complex{4, 0}}},
+      // Random value test
+      Array{Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}},
+            Marray{Complex{0.5, 0.5}, Complex{1.2, 1.2}, Complex{-2.8, -2.8},
+                   Complex{3.7, 3.7}}},
+      // Repeated value test
+      Array{Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}},
+            Marray{Complex{1, 1}, Complex{1, 1}, Complex{1, 1}, Complex{1, 1}}},
+      // Negative value test
+      Array{Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}},
+            Marray{Complex{-3.0, -3.0}, Complex{2.5, 2.5}, Complex{-1.2, -1.2},
+                   Complex{0, 0}}},
+      // Large value test
+      Array{
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}},
+          Marray{Complex{1000000.0, 1000000.0}, Complex{2000000.0, 2000000.0},
+                 Complex{3000000.0, 3000000.0}, Complex{4000000.0, 4000000.0}}},
+      // Small value test
+      Array{Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}},
+            Marray{Complex{0.0001, 0.0001}, Complex{0.0002, 0.0002},
+                   Complex{0.0003, 0.0003}, Complex{0.0004, 0.0004}}},
+      // Edge case value test
+      Array{
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{
+              Complex{nan_val<T>, nan_val<T>}, Complex{inf_val<T>, inf_val<T>},
+              Complex{nan_val<T>, inf_val<T>}, Complex{inf_val<T>, nan_val<T>}},
+          Marray{Complex{nan_val<T>, nan_val<T>},
+                 Complex{inf_val<T>, inf_val<T>},
+                 Complex{nan_val<T>, inf_val<T>},
+                 Complex{inf_val<T>, nan_val<T>}}});
+  const auto binary_op = BinaryOperation{};
+
+  if (is_type_supported<T>(q)) {
+    test_reduce_over_group(q, test_cases, binary_op);
+  }
+}
diff --git a/tests/test_conversion.cpp b/tests/test_conversion.cpp
new file mode 100644
index 0000000..2e5f19e
--- /dev/null
+++ b/tests/test_conversion.cpp
@@ -0,0 +1,70 @@
+#include "test_helper.hpp"
+
+#include <array>
+#include <cmath>
+
+// Check conversion sycl:complex to std::complex
+TEMPLATE_TEST_CASE("Test sycl::complex to std::complex conversion",
+                   "[conversion]", double, float, sycl::half) {
+  using T = TestType;
+
+  auto arr = std::array<sycl::ext::cplx::complex<T>, 8>{
+      sycl::ext::cplx::complex<T>{0, 0},  sycl::ext::cplx::complex<T>{-0, 0},
+      sycl::ext::cplx::complex<T>{0, -0}, sycl::ext::cplx::complex<T>{-0, -0},
+
+      sycl::ext::cplx::complex<T>{1, 1},  sycl::ext::cplx::complex<T>{-1, 1},
+      sycl::ext::cplx::complex<T>{1, -1}, sycl::ext::cplx::complex<T>{-1, -1},
+  };
+
+  for (const auto &lhs : arr) {
+    const auto rhs = static_cast<std::complex<T>>(lhs);
+
+    assert(lhs.real() == rhs.real() &&
+           "sycl::complex differs from std::complex after conversion");
+    assert(lhs.imag() == rhs.imag() &&
+           "sycl::complex differs from std::complex after conversion");
+  }
+}
+
+// Check edge-cases conversion sycl:complex to std::complex
+TEMPLATE_TEST_CASE("Test edge-cases sycl::complex to std::complex conversion",
+                   "[conversion]", double, float, sycl::half) {
+  using T = TestType;
+
+  {
+    auto lhs = sycl::ext::cplx::complex<T>{inf_val<T>, inf_val<T>};
+    auto rhs = static_cast<std::complex<T>>(lhs);
+
+    assert(std::isinf(lhs.real()) && std::isinf(rhs.real()) &&
+           "sycl::complex differs from std::complex after conversion");
+    assert(std::isinf(lhs.imag()) && std::isinf(rhs.imag()) &&
+           "sycl::complex differs from std::complex after conversion");
+  }
+  {
+    auto lhs = sycl::ext::cplx::complex<T>{inf_val<T>, nan_val<T>};
+    auto rhs = static_cast<std::complex<T>>(lhs);
+
+    assert(std::isinf(lhs.real()) && std::isinf(rhs.real()) &&
+           "sycl::complex differs from std::complex after conversion");
+    assert(std::isnan(lhs.imag()) && std::isnan(rhs.imag()) &&
+           "sycl::complex differs from std::complex after conversion");
+  }
+  {
+    auto lhs = sycl::ext::cplx::complex<T>{nan_val<T>, inf_val<T>};
+    auto rhs = static_cast<std::complex<T>>(lhs);
+
+    assert(std::isnan(lhs.real()) && std::isnan(rhs.real()) &&
+           "sycl::complex differs from std::complex after conversion");
+    assert(std::isinf(lhs.imag()) && std::isinf(rhs.imag()) &&
+           "sycl::complex differs from std::complex after conversion");
+  }
+  {
+    auto lhs = sycl::ext::cplx::complex<T>{nan_val<T>, nan_val<T>};
+    auto rhs = static_cast<std::complex<T>>(lhs);
+
+    assert(std::isnan(lhs.real()) && std::isnan(rhs.real()) &&
+           "sycl::complex differs from std::complex after conversion");
+    assert(std::isnan(lhs.imag()) && std::isnan(rhs.imag()) &&
+           "sycl::complex differs from std::complex after conversion");
+  }
+}
diff --git a/tests/test_gencomplex.cpp b/tests/test_gencomplex.cpp
deleted file mode 100644
index 64b0bd6..0000000
--- a/tests/test_gencomplex.cpp
+++ /dev/null
@@ -1,17 +0,0 @@
-#include "test_helper.hpp"
-
-using namespace sycl::ext::cplx;
-
-// Check is_gencomplex
-TEST_CASE("Test is_gencomplex", "[gencomplex]") {
-  static_assert(is_gencomplex<complex<double>>::value == true);
-  static_assert(is_gencomplex<complex<float>>::value == true);
-  static_assert(is_gencomplex<complex<sycl::half>>::value == true);
-
-  static_assert(is_gencomplex<complex<long long>>::value == false);
-  static_assert(is_gencomplex<complex<long>>::value == false);
-  static_assert(is_gencomplex<complex<int>>::value == false);
-  static_assert(is_gencomplex<complex<unsigned long long>>::value == false);
-  static_assert(is_gencomplex<complex<unsigned long>>::value == false);
-  static_assert(is_gencomplex<complex<unsigned int>>::value == false);
-}
diff --git a/tests/test_helper.hpp b/tests/test_helper.hpp
index 773e980..a5a2ec4 100644
--- a/tests/test_helper.hpp
+++ b/tests/test_helper.hpp
@@ -1,8 +1,8 @@
 #include <catch2/catch_template_test_macros.hpp>
 #include <catch2/catch_test_macros.hpp>
 #include <catch2/generators/catch_generators.hpp>
-#include <cmath>
-#include <iomanip>
+#include <catch2/matchers/catch_matchers_floating_point.hpp>
+
 #include <sycl/sycl.hpp>
 
 #include "sycl_ext_complex.hpp"
@@ -67,6 +67,8 @@ template <> struct numeric_limits<sycl::half> {
 } // namespace std
 #endif
 
+namespace detail {
+
 template <typename T, typename U> inline bool is_nan_or_inf(T x, U y) {
   return (std::isnan(x.real()) && std::isnan(y.real())) ||
          (std::isnan(x.imag()) && std::isnan(y.imag())) ||
@@ -74,7 +76,6 @@ template <typename T, typename U> inline bool is_nan_or_inf(T x, U y) {
          (std::isinf(x.imag()) && std::isinf(y.imag()));
 }
 
-namespace detail {
 template <typename T> bool almost_equal(T x, T y, int ulp) {
   if (std::isnan(x) && std::isnan(y))
     return true;
@@ -87,11 +88,12 @@ template <typename T> bool almost_equal(T x, T y, int ulp) {
 }
 
 template <typename T>
-bool almost_equal(sycl::ext::cplx::complex<T> x, std::complex<T> y, int ulp) {
-  auto diff = std::abs((std::complex<T>)x - y);
+bool almost_equal(std::complex<T> output, std::complex<T> reference, int ulp) {
+  auto diff = std::abs(output - reference);
   return diff <= std::numeric_limits<T>::epsilon() *
-                     std::abs((std::complex<T>)x + y) * ulp ||
-         diff < std::numeric_limits<T>::min() || is_nan_or_inf(x, y);
+                     std::abs(output + reference) * ulp ||
+         diff < std::numeric_limits<T>::min() ||
+         is_nan_or_inf(output, reference);
 }
 
 // Helpers for testing half
@@ -106,6 +108,7 @@ inline std::complex<float> trunc_float(std::complex<float> c) {
   auto c_sycl_half = static_cast<sycl::ext::cplx::complex<sycl::half>>(c);
   return sycl_half_to_float(c_sycl_half);
 }
+
 } // namespace detail
 
 // Helper for initializing std::complex values for tests only needed because
@@ -164,34 +167,68 @@ auto constexpr convert_marray(sycl::marray<std::complex<Tin>, NumElements> c) {
 
 // Helpers for comparing SyclCPLX and standard c++ results
 
+namespace detail {
+
+template <typename T>
+struct is_a_complex
+    : std::integral_constant<bool,
+                             sycl::ext::cplx::is_gencomplex_v<T> ||
+                                 std::is_same_v<std::complex<sycl::half>, T> ||
+                                 std::is_same_v<std::complex<float>, T> ||
+                                 std::is_same_v<std::complex<double>, T>> {};
+template <typename T>
+inline constexpr bool is_a_complex_v = is_a_complex<T>::value;
+
+} // namespace detail
+
+/// Specialization for double, float, sycl::half
 template <typename T>
-void check_results(sycl::ext::cplx::complex<T> output,
-                   std::complex<T> reference, int tol_multiplier = 1) {
+typename std::enable_if_t<sycl::ext::cplx::is_genfloat_v<T>, void>
+check_results(T output, T reference, int tol_multiplier = 1) {
   CHECK(detail::almost_equal(output, reference,
                              tol_multiplier * SYCL_CPLX_TOL_ULP));
 }
 
-template <typename T>
-void check_results(T output, T reference, int tol_multiplier = 1) {
-  CHECK(detail::almost_equal(output, reference,
+/// Specialization for sycl::complex and std::complex
+template <typename LHS, typename RHS>
+typename std::enable_if_t<
+    detail::is_a_complex_v<LHS> && detail::is_a_complex_v<RHS>, void>
+check_results(LHS output, RHS reference, int tol_multiplier = 1) {
+  using T1 = typename LHS::value_type;
+  using T2 = typename RHS::value_type;
+
+  if (!std::is_same_v<T1, T2>) {
+    FAIL("check_results can be called with sycl::complex and/or std::complex "
+         "but with the same value_type");
+  }
+
+  CHECK(detail::almost_equal(static_cast<std::complex<T1>>(output),
+                             static_cast<std::complex<T2>>(reference),
                              tol_multiplier * SYCL_CPLX_TOL_ULP));
 }
 
-template <typename T, std::size_t NumElements>
-void check_results(
-    sycl::marray<sycl::ext::cplx::complex<T>, NumElements> output,
-    sycl::marray<std::complex<T>, NumElements> reference,
-    int tol_multiplier = 1) {
-  for (std::size_t i = 0; i < NumElements; ++i) {
+/// Specialization for sycl::marray
+template <typename LHS, std::size_t NumElementsLHS, typename RHS,
+          std::size_t NumElementsRHS>
+void check_results(sycl::marray<LHS, NumElementsLHS> output,
+                   sycl::marray<RHS, NumElementsRHS> reference,
+                   int tol_multiplier = 1) {
+  if (NumElementsLHS != NumElementsRHS) {
+    FAIL("check_results can be called with sycl::marray but with the same "
+         "NumElement");
+  }
+
+  for (std::size_t i = 0; i < NumElementsLHS; ++i) {
     check_results(output[i], reference[i], tol_multiplier);
   }
 }
 
+/// Specialization for std::array
 template <typename T, std::size_t NumElements>
-void check_results(sycl::marray<T, NumElements> output,
-                   sycl::marray<T, NumElements> reference,
+void check_results(std::array<T, NumElements> output,
+                   std::array<T, NumElements> reference,
                    int tol_multiplier = 1) {
   for (std::size_t i = 0; i < NumElements; ++i) {
-    check_results<T>(output[i], reference[i], tol_multiplier);
+    check_results(output[i], reference[i], tol_multiplier);
   }
 }
diff --git a/tests/test_traits.cpp b/tests/test_traits.cpp
new file mode 100644
index 0000000..7473862
--- /dev/null
+++ b/tests/test_traits.cpp
@@ -0,0 +1,130 @@
+#include "test_helper.hpp"
+
+using namespace sycl::ext::cplx;
+
+// Check is_gencomplex
+TEST_CASE("Test is_gencomplex", "[traits]") {
+  static_assert(is_gencomplex<complex<double>>::value == true);
+  static_assert(is_gencomplex<complex<float>>::value == true);
+  static_assert(is_gencomplex<complex<sycl::half>>::value == true);
+
+  static_assert(is_gencomplex_v<complex<double>> == true);
+  static_assert(is_gencomplex_v<complex<float>> == true);
+  static_assert(is_gencomplex_v<complex<sycl::half>> == true);
+
+  static_assert(is_gencomplex<complex<long long>>::value == false);
+  static_assert(is_gencomplex<complex<long>>::value == false);
+  static_assert(is_gencomplex<complex<int>>::value == false);
+  static_assert(is_gencomplex<complex<unsigned long long>>::value == false);
+  static_assert(is_gencomplex<complex<unsigned long>>::value == false);
+  static_assert(is_gencomplex<complex<unsigned int>>::value == false);
+
+  static_assert(is_gencomplex_v<complex<long long>> == false);
+  static_assert(is_gencomplex_v<complex<long>> == false);
+  static_assert(is_gencomplex_v<complex<int>> == false);
+  static_assert(is_gencomplex_v<complex<unsigned long long>> == false);
+  static_assert(is_gencomplex_v<complex<unsigned long>> == false);
+  static_assert(is_gencomplex_v<complex<unsigned int>> == false);
+}
+
+// Check is_genfloat
+TEST_CASE("Test is_genfloat", "[traits]") {
+  static_assert(is_genfloat<double>::value == true);
+  static_assert(is_genfloat<float>::value == true);
+  static_assert(is_genfloat<sycl::half>::value == true);
+
+  static_assert(is_genfloat_v<double> == true);
+  static_assert(is_genfloat_v<float> == true);
+  static_assert(is_genfloat_v<sycl::half> == true);
+
+  static_assert(is_genfloat<long long>::value == false);
+  static_assert(is_genfloat<long>::value == false);
+  static_assert(is_genfloat<int>::value == false);
+  static_assert(is_genfloat<unsigned long long>::value == false);
+  static_assert(is_genfloat<unsigned long>::value == false);
+  static_assert(is_genfloat<unsigned int>::value == false);
+
+  static_assert(is_genfloat_v<long long> == false);
+  static_assert(is_genfloat_v<long> == false);
+  static_assert(is_genfloat_v<int> == false);
+  static_assert(is_genfloat_v<unsigned long long> == false);
+  static_assert(is_genfloat_v<unsigned long> == false);
+  static_assert(is_genfloat_v<unsigned int> == false);
+}
+
+// Check is_mgencomplex
+TEST_CASE("Test is_mgencomplex", "[traits]") {
+  static_assert(is_mgencomplex<sycl::marray<complex<double>, 42>>::value ==
+                true);
+  static_assert(is_mgencomplex<sycl::marray<complex<float>, 42>>::value ==
+                true);
+  static_assert(is_mgencomplex<sycl::marray<complex<sycl::half>, 42>>::value ==
+                true);
+
+  static_assert(is_mgencomplex_v<sycl::marray<complex<double>, 42>> == true);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<float>, 42>> == true);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<sycl::half>, 42>> ==
+                true);
+
+  static_assert(is_mgencomplex<sycl::marray<complex<long long>, 42>>::value ==
+                false);
+  static_assert(is_mgencomplex<sycl::marray<complex<long>, 42>>::value ==
+                false);
+  static_assert(is_mgencomplex<sycl::marray<complex<int>, 42>>::value == false);
+  static_assert(
+      is_mgencomplex<sycl::marray<complex<unsigned long long>, 42>>::value ==
+      false);
+  static_assert(
+      is_mgencomplex<sycl::marray<complex<unsigned long>, 42>>::value == false);
+  static_assert(
+      is_mgencomplex<sycl::marray<complex<unsigned int>, 42>>::value == false);
+
+  static_assert(is_mgencomplex_v<sycl::marray<complex<long long>, 42>> ==
+                false);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<long>, 42>> == false);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<int>, 42>> == false);
+  static_assert(
+      is_mgencomplex_v<sycl::marray<complex<unsigned long long>, 42>> == false);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<unsigned long>, 42>> ==
+                false);
+  static_assert(is_mgencomplex_v<sycl::marray<complex<unsigned int>, 42>> ==
+                false);
+}
+
+// Check is_plus
+TEST_CASE("Test is_plus", "[traits]") {
+  static_assert(cplex::detail::is_plus_v<sycl::plus<>> == true);
+
+  static_assert(cplex::detail::is_plus_v<sycl::multiplies<>> == false);
+  static_assert(cplex::detail::is_plus_v<sycl::bit_and<>> == false);
+  static_assert(cplex::detail::is_plus_v<sycl::logical_and<>> == false);
+  static_assert(cplex::detail::is_plus_v<sycl::minimum<>> == false);
+  static_assert(cplex::detail::is_plus_v<sycl::maximum<>> == false);
+}
+
+// Check is_multiplies
+TEST_CASE("Test is_multiplies", "[traits]") {
+  static_assert(cplex::detail::is_multiplies_v<sycl::multiplies<>> == true);
+
+  static_assert(cplex::detail::is_multiplies_v<sycl::plus<>> == false);
+  static_assert(cplex::detail::is_multiplies_v<sycl::bit_and<>> == false);
+  static_assert(cplex::detail::is_multiplies_v<sycl::logical_and<>> == false);
+  static_assert(cplex::detail::is_multiplies_v<sycl::minimum<>> == false);
+  static_assert(cplex::detail::is_multiplies_v<sycl::maximum<>> == false);
+}
+
+// Check is_binary_op_supported
+TEST_CASE("Test is_binary_op_supported", "[traits]") {
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::plus<>> == true);
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::multiplies<>> ==
+                true);
+
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::bit_and<>> ==
+                false);
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::logical_and<>> ==
+                false);
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::minimum<>> ==
+                false);
+  static_assert(cplex::detail::is_binary_op_supported_v<sycl::maximum<>> ==
+                false);
+}