diff --git a/sycl/include/sycl/group_algorithm.hpp b/sycl/include/sycl/group_algorithm.hpp
index 8963184814b4..5125ceca5fc0 100644
--- a/sycl/include/sycl/group_algorithm.hpp
+++ b/sycl/include/sycl/group_algorithm.hpp
@@ -143,6 +143,19 @@ using is_plus_or_multiplies_if_complex = std::integral_constant<
                                    is_multiplies<T, BinaryOperation>::value)
                                 : std::true_type::value)>;
 
+// used to transform a vector op to a scalar op;
+// e.g. sycl::plus<std::vec<T, N>> to sycl::plus<T>
+template <typename T> struct get_scalar_binary_op;
+
+template <template <typename> typename F, typename T, int n>
+struct get_scalar_binary_op<F<sycl::vec<T, n>>> {
+  using type = F<T>;
+};
+
+template <template <typename> typename F> struct get_scalar_binary_op<F<void>> {
+  using type = F<void>;
+};
+
 // ---- identity_for_ga_op
 //   the group algorithms support std::complex, limited to sycl::plus operation
 //   get the correct identity for group algorithm operation.
@@ -200,11 +213,8 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 reduce_over_group(Group g, T x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same_v<decltype(binary_op(x, x)), T> ||
-          (std::is_same_v<T, half> &&
-           std::is_same_v<decltype(binary_op(x, x)), float>),
+      std::is_same_v<decltype(binary_op(x, x)), T>,
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<__spv::GroupOperation::Reduce>(
@@ -239,24 +249,21 @@ reduce_over_group(Group g, T x, BinaryOperation binary_op) {
 #endif
 }
 
-template <typename Group, typename T, int N, class BinaryOperation>
-std::enable_if_t<
-    (is_group_v<std::decay_t<Group>> &&
-     detail::is_vector_arithmetic_or_complex<sycl::vec<T, N>>::value &&
-     detail::is_native_op<sycl::vec<T, N>, BinaryOperation>::value),
-    sycl::vec<T, N>>
-reduce_over_group(Group g, sycl::vec<T, N> x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
+template <typename Group, typename T, class BinaryOperation>
+std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
+                  detail::is_vector_arithmetic_or_complex<T>::value &&
+                  detail::is_native_op<T, BinaryOperation>::value),
+                 T>
+reduce_over_group(Group g, T x, BinaryOperation binary_op) {
   static_assert(
-      std::is_same_v<decltype(binary_op(x[0], x[0])),
-                     typename sycl::vec<T, N>::element_type> ||
-          (std::is_same_v<sycl::vec<T, N>, half> &&
-           std::is_same_v<decltype(binary_op(x[0], x[0])), float>),
+      std::is_same_v<decltype(binary_op(x, x)), T>,
       "Result type of binary_op must match reduction accumulation type.");
-  sycl::vec<T, N> result;
-
-  detail::loop<N>(
-      [&](size_t s) { result[s] = reduce_over_group(g, x[s], binary_op); });
+  T result;
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
+  detail::loop<x.size()>([&](size_t s) {
+    result[s] = reduce_over_group(g, x[s], scalar_binary_op);
+  });
   return result;
 }
 
@@ -272,11 +279,8 @@ std::enable_if_t<
      std::is_convertible_v<V, T>),
     T>
 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same_v<decltype(binary_op(init, x)), T> ||
-          (std::is_same_v<T, half> &&
-           std::is_same_v<decltype(binary_op(init, x)), float>),
+      std::is_same_v<decltype(binary_op(init, x)), T>,
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return binary_op(init, reduce_over_group(g, T(x), binary_op));
@@ -295,17 +299,16 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 reduce_over_group(Group g, V x, T init, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same_v<decltype(binary_op(init[0], x[0])),
-                     typename T::element_type> ||
-          (std::is_same_v<T, half> &&
-           std::is_same_v<decltype(binary_op(init[0], x[0])), float>),
+      std::is_same_v<decltype(binary_op(init, x)), T>,
       "Result type of binary_op must match reduction accumulation type.");
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
 #ifdef __SYCL_DEVICE_ONLY__
   T result = init;
   for (int s = 0; s < x.size(); ++s) {
-    result[s] = binary_op(init[s], reduce_over_group(g, x[s], binary_op));
+    result[s] =
+        scalar_binary_op(init[s], reduce_over_group(g, x[s], scalar_binary_op));
   }
   return result;
 #else
@@ -326,11 +329,8 @@ std::enable_if_t<
      detail::is_native_op<T, BinaryOperation>::value),
     T>
 joint_reduce(Group g, Ptr first, Ptr last, T init, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
   static_assert(
-      std::is_same_v<decltype(binary_op(init, *first)), T> ||
-          (std::is_same_v<T, half> &&
-           std::is_same_v<decltype(binary_op(init, *first)), float>),
+      std::is_same_v<decltype(binary_op(init, *first)), T>,
       "Result type of binary_op must match reduction accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T partial = detail::identity_for_ga_op<T, BinaryOperation>();
@@ -630,10 +630,7 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(x, x)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(x, x)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<__spv::GroupOperation::ExclusiveScan>(
@@ -674,15 +671,13 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 exclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(x[0], x[0])),
-                               typename T::element_type> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(x[0], x[0])), float>),
+  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
   T result;
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
   for (int s = 0; s < x.size(); ++s) {
-    result[s] = exclusive_scan_over_group(g, x[s], binary_op);
+    result[s] = exclusive_scan_over_group(g, x[s], scalar_binary_op);
   }
   return result;
 }
@@ -697,15 +692,13 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init[0], x[0])),
-                               typename T::element_type> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init[0], x[0])), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
   T result;
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
   for (int s = 0; s < x.size(); ++s) {
-    result[s] = exclusive_scan_over_group(g, x[s], init[s], binary_op);
+    result[s] = exclusive_scan_over_group(g, x[s], init[s], scalar_binary_op);
   }
   return result;
 }
@@ -720,10 +713,7 @@ std::enable_if_t<
      std::is_convertible_v<V, T>),
     T>
 exclusive_scan_over_group(Group g, V x, T init, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init, x)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init, x)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   typename Group::linear_id_type local_linear_id =
@@ -760,10 +750,7 @@ std::enable_if_t<
     OutPtr>
 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result, T init,
                      BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init, *first)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init, *first)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
@@ -815,14 +802,9 @@ std::enable_if_t<
     OutPtr>
 joint_exclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                      BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(
-      std::is_same_v<decltype(binary_op(*first, *first)),
-                     typename detail::remove_pointer<OutPtr>::type> ||
-          (std::is_same_v<typename detail::remove_pointer<OutPtr>::type,
-                          half> &&
-           std::is_same_v<decltype(binary_op(*first, *first)), float>),
-      "Result type of binary_op must match scan accumulation type.");
+  static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
+                               typename detail::remove_pointer<OutPtr>::type>,
+                "Result type of binary_op must match scan accumulation type.");
   using T = typename detail::remove_pointer<OutPtr>::type;
   T init = detail::identity_for_ga_op<T, BinaryOperation>();
   return joint_exclusive_scan(g, first, last, result, init, binary_op);
@@ -838,15 +820,13 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(x[0], x[0])),
-                               typename T::element_type> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(x[0], x[0])), float>),
+  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
   T result;
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
   for (int s = 0; s < x.size(); ++s) {
-    result[s] = inclusive_scan_over_group(g, x[s], binary_op);
+    result[s] = inclusive_scan_over_group(g, x[s], scalar_binary_op);
   }
   return result;
 }
@@ -859,10 +839,7 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 inclusive_scan_over_group(Group g, T x, BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(x, x)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(x, x)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(x, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   return sycl::detail::calc<__spv::GroupOperation::InclusiveScan>(
@@ -908,10 +885,7 @@ std::enable_if_t<
      std::is_convertible_v<V, T>),
     T>
 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init, x)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init, x)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   T y = x;
@@ -934,14 +908,13 @@ std::enable_if_t<(is_group_v<std::decay_t<Group>> &&
                   detail::is_native_op<T, BinaryOperation>::value),
                  T>
 inclusive_scan_over_group(Group g, V x, BinaryOperation binary_op, T init) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init[0], x[0])), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init[0], x[0])), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, x)), T>,
                 "Result type of binary_op must match scan accumulation type.");
   T result;
+  typename detail::get_scalar_binary_op<BinaryOperation>::type
+      scalar_binary_op{};
   for (int s = 0; s < x.size(); ++s) {
-    result[s] = inclusive_scan_over_group(g, x[s], binary_op, init[s]);
+    result[s] = inclusive_scan_over_group(g, x[s], scalar_binary_op, init[s]);
   }
   return result;
 }
@@ -962,10 +935,7 @@ std::enable_if_t<
     OutPtr>
 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                      BinaryOperation binary_op, T init) {
-  // FIXME: Do not special-case for half precision
-  static_assert(std::is_same_v<decltype(binary_op(init, *first)), T> ||
-                    (std::is_same_v<T, half> &&
-                     std::is_same_v<decltype(binary_op(init, *first)), float>),
+  static_assert(std::is_same_v<decltype(binary_op(init, *first)), T>,
                 "Result type of binary_op must match scan accumulation type.");
 #ifdef __SYCL_DEVICE_ONLY__
   ptrdiff_t offset = sycl::detail::get_local_linear_id(g);
@@ -1014,14 +984,9 @@ std::enable_if_t<
     OutPtr>
 joint_inclusive_scan(Group g, InPtr first, InPtr last, OutPtr result,
                      BinaryOperation binary_op) {
-  // FIXME: Do not special-case for half precision
-  static_assert(
-      std::is_same_v<decltype(binary_op(*first, *first)),
-                     typename detail::remove_pointer<OutPtr>::type> ||
-          (std::is_same_v<typename detail::remove_pointer<OutPtr>::type,
-                          half> &&
-           std::is_same_v<decltype(binary_op(*first, *first)), float>),
-      "Result type of binary_op must match scan accumulation type.");
+  static_assert(std::is_same_v<decltype(binary_op(*first, *first)),
+                               typename detail::remove_pointer<OutPtr>::type>,
+                "Result type of binary_op must match scan accumulation type.");
 
   using T = typename detail::remove_pointer<OutPtr>::type;
   T init = detail::identity_for_ga_op<T, BinaryOperation>();
diff --git a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp
index 152b264b123c..a836cab08684 100644
--- a/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/exclusive_scan_sycl2020.cpp
@@ -33,18 +33,20 @@ OutputIterator exclusive_scan(InputIterator first, InputIterator last,
 }
 } // namespace emu
 
+queue q;
+
 template <typename SpecializationKernelName, typename InputContainer,
-          typename OutputContainer, class BinaryOperation>
-void test(queue q, InputContainer input, OutputContainer output,
-          BinaryOperation binary_op,
-          typename OutputContainer::value_type identity) {
+          class BinaryOperation>
+void test(const InputContainer &input, BinaryOperation binary_op,
+          typename InputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
-  typedef typename OutputContainer::value_type OutputT;
+  typedef InputT OutputT;
   typedef class exclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
-  constexpr size_t G = 64;
-  constexpr size_t N = input.size(); // 128 or 12
+  constexpr size_t G = 16;
+  constexpr size_t N = std::tuple_size_v<InputContainer>; // 128 or 12
   constexpr size_t confirmRange = std::min(G, N);
-  std::vector<OutputT> expected(N);
+  std::array<OutputT, N> output;
+  std::array<OutputT, N> expected;
 
   // checking
   // template <typename Group, typename T, typename BinaryOperation>
@@ -64,11 +66,11 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
                       expected.begin(), identity, binary_op);
-  assert(std::equal(output.begin(), output.begin() + confirmRange,
-                    expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + confirmRange,
+                      expected.begin()));
 
   typedef class exclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
-  constexpr OutputT init = 42;
+  constexpr OutputT init(42);
 
   // checking
   // template <typename Group, typename V, typename T, class BinaryOperation>
@@ -89,8 +91,8 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::exclusive_scan(input.begin(), input.begin() + confirmRange,
                       expected.begin(), init, binary_op);
-  assert(std::equal(output.begin(), output.begin() + confirmRange,
-                    expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + confirmRange,
+                      expected.begin()));
 
   typedef class exclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
 
@@ -117,7 +119,7 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(),
                       identity, binary_op);
-  assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + N, expected.begin()));
 
   typedef class exclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
 
@@ -145,7 +147,7 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::exclusive_scan(input.begin(), input.begin() + N, expected.begin(), init,
                       binary_op);
-  assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + N, expected.begin()));
 }
 
 int main() {
@@ -157,50 +159,51 @@ int main() {
 
   constexpr int N = 128;
   std::array<int, N> input;
-  std::array<int, N> output;
   std::iota(input.begin(), input.end(), 2);
-  std::fill(output.begin(), output.end(), 0);
 
   // Smaller size as the multiplication test
   // will result in computing of a factorial
   // 12! fits in a 32 bits integer.
   constexpr int M = 12;
   std::array<int, M> input_small;
-  std::array<int, M> output_small;
   std::iota(input_small.begin(), input_small.end(), 1);
-  std::fill(output_small.begin(), output_small.end(), 0);
 
-  test<class KernelNamePlusV>(q, input, output, sycl::plus<>(), 0);
-  test<class KernelNameMinimumV>(q, input, output, sycl::minimum<>(),
+  test<class KernelNamePlusV>(input, sycl::plus<>(), 0);
+  test<class KernelNameMinimumV>(input, sycl::minimum<>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumV>(q, input, output, sycl::maximum<>(),
+  test<class KernelNameMaximumV>(input, sycl::maximum<>(),
                                  std::numeric_limits<int>::lowest());
 
-  test<class KernelNamePlusI>(q, input, output, sycl::plus<int>(), 0);
-  test<class KernelNameMinimumI>(q, input, output, sycl::minimum<int>(),
+  test<class KernelNamePlusI>(input, sycl::plus<int>(), 0);
+  test<class KernelNameMinimumI>(input, sycl::minimum<int>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
+  test<class KernelNameMaximumI>(input, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-  test<class KernelNameMultipliesI>(q, input_small, output_small,
-                                    sycl::multiplies<int>(), 1);
-  test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
-  test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
-  test<class KernelNameBitAndI>(q, input_small, output_small,
-                                sycl::bit_and<int>(), ~0);
+  test<class KernelNameMultipliesI>(input_small, sycl::multiplies<int>(), 1);
+  test<class KernelNameBitOrI>(input, sycl::bit_or<int>(), 0);
+  test<class KernelNameBitXorI>(input, sycl::bit_xor<int>(), 0);
+  test<class KernelNameBitAndI>(input_small, sycl::bit_and<int>(), ~0);
 
-  test<class LogicalOrInt>(q, input, output, sycl::logical_or<int>(), 0);
-  test<class LogicalAndInt>(q, input, output, sycl::logical_and<int>(), 1);
+  test<class LogicalOrInt>(input, sycl::logical_or<int>(), 0);
+  test<class LogicalAndInt>(input, sycl::logical_and<int>(), 1);
 
   std::array<bool, N> bool_input = {};
-  std::array<bool, N> bool_output = {};
-  test<class LogicalOrBool>(q, bool_input, bool_output,
-                            sycl::logical_or<bool>(), false);
-  test<class LogicalOrVoid>(q, bool_input, bool_output, sycl::logical_or<>(),
-                            false);
-  test<class LogicalAndBool>(q, bool_input, bool_output,
-                             sycl::logical_and<bool>(), true);
-  test<class LogicalAndVoid>(q, bool_input, bool_output, sycl::logical_and<>(),
-                             true);
+  test<class LogicalOrBool>(bool_input, sycl::logical_or<bool>(), false);
+  test<class LogicalOrVoid>(bool_input, sycl::logical_or<>(), false);
+  test<class LogicalAndBool>(bool_input, sycl::logical_and<bool>(), true);
+  test<class LogicalAndVoid>(bool_input, sycl::logical_and<>(), true);
+
+  std::array<int2, N> int2_input = {};
+  std::iota(int2_input.begin(), int2_input.end(), 0);
+  test<class PlusInt2>(int2_input, sycl::plus<int2>(), {0, 0});
+  test<class PlusInt2V>(int2_input, sycl::plus<>(), {0, 0});
+
+  if (q.get_device().has(aspect::fp16)) {
+    std::array<half, 32> half_input = {};
+    std::iota(half_input.begin(), half_input.end(), 0);
+    test<class PlusHalf>(half_input, sycl::plus<half>(), 0);
+    test<class PlusHalfV>(half_input, sycl::plus<>(), 0);
+  }
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_exclusive_scan and
@@ -208,22 +211,17 @@ int main() {
   // sycl::plus binary operation.
 #ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
   std::array<std::complex<float>, N> input_cf;
-  std::array<std::complex<float>, N> output_cf;
   std::iota(input_cf.begin(), input_cf.end(), 0);
-  std::fill(output_cf.begin(), output_cf.end(), 0);
-  test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
+  test<class KernelNamePlusComplexF>(input_cf,
                                      sycl::plus<std::complex<float>>(), 0);
-  test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
+  test<class KernelNamePlusUnspecF>(input_cf, sycl::plus<>(), 0);
 
   if (q.get_device().has(aspect::fp64)) {
     std::array<std::complex<double>, N> input_cd;
-    std::array<std::complex<double>, N> output_cd;
     std::iota(input_cd.begin(), input_cd.end(), 0);
-    std::fill(output_cd.begin(), output_cd.end(), 0);
-    test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
+    test<class KernelNamePlusComplexD>(input_cd,
                                        sycl::plus<std::complex<double>>(), 0);
-    test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
-                                      0);
+    test<class KernelNamePlusUnspecD>(input_cd, sycl::plus<>(), 0);
   } else {
     std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
               << std::endl;
diff --git a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp
index d3eef274ee25..8ff5da727260 100644
--- a/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/inclusive_scan_sycl2020.cpp
@@ -12,6 +12,8 @@
 #include <vector>
 using namespace sycl;
 
+queue q;
+
 template <class SpecializationKernelName, int TestNumber>
 class inclusive_scan_kernel;
 
@@ -34,17 +36,17 @@ OutputIterator inclusive_scan(InputIterator first, InputIterator last,
 } // namespace emu
 
 template <typename SpecializationKernelName, typename InputContainer,
-          typename OutputContainer, class BinaryOperation>
-void test(queue q, InputContainer input, OutputContainer output,
-          BinaryOperation binary_op,
-          typename OutputContainer::value_type identity) {
+          class BinaryOperation>
+void test(const InputContainer &input, BinaryOperation binary_op,
+          typename InputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
-  typedef typename OutputContainer::value_type OutputT;
+  typedef InputT OutputT;
   typedef class inclusive_scan_kernel<SpecializationKernelName, 0> kernel_name0;
-  constexpr size_t N = input.size(); // 128 or 12
-  constexpr size_t G = 64;
+  constexpr size_t N = std::tuple_size_v<InputContainer>; // 128 or 12
+  constexpr size_t G = 16;
   constexpr size_t confirmRange = std::min(G, N);
-  std::vector<OutputT> expected(N);
+  std::array<OutputT, N> output;
+  std::array<OutputT, N> expected;
 
   // checking
   // template <typename Group, typename T, class BinaryOperation>
@@ -64,11 +66,11 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
                       expected.begin(), binary_op, identity);
-  assert(std::equal(output.begin(), output.begin() + confirmRange,
-                    expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + confirmRange,
+                      expected.begin()));
 
   typedef class inclusive_scan_kernel<SpecializationKernelName, 1> kernel_name1;
-  constexpr OutputT init = 42;
+  constexpr OutputT init(42);
 
   // checking
   // template <typename Group, typename V, class BinaryOperation, typename T>
@@ -89,8 +91,8 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::inclusive_scan(input.begin(), input.begin() + confirmRange,
                       expected.begin(), binary_op, init);
-  assert(std::equal(output.begin(), output.begin() + confirmRange,
-                    expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + confirmRange,
+                      expected.begin()));
 
   typedef class inclusive_scan_kernel<SpecializationKernelName, 2> kernel_name2;
 
@@ -116,7 +118,7 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(),
                       binary_op, identity);
-  assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + N, expected.begin()));
 
   typedef class inclusive_scan_kernel<SpecializationKernelName, 3> kernel_name3;
 
@@ -143,7 +145,7 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   emu::inclusive_scan(input.begin(), input.begin() + N, expected.begin(),
                       binary_op, init);
-  assert(std::equal(output.begin(), output.begin() + N, expected.begin()));
+  assert(ranges_equal(output.begin(), output.begin() + N, expected.begin()));
 }
 
 int main() {
@@ -155,50 +157,56 @@ int main() {
 
   constexpr int N = 128;
   std::array<int, N> input;
-  std::array<int, N> output;
   std::iota(input.begin(), input.end(), 2);
-  std::fill(output.begin(), output.end(), 0);
 
   // Smaller size as the multiplication test
   // will result in computing of a factorial
   // 12! fits in a 32 bits integer.
   constexpr int M = 12;
   std::array<int, M> input_small;
-  std::array<int, M> output_small;
   std::iota(input_small.begin(), input_small.end(), 1);
-  std::fill(output_small.begin(), output_small.end(), 0);
 
-  test<class KernelNamePlusV>(q, input, output, sycl::plus<>(), 0);
-  test<class KernelNameMinimumV>(q, input, output, sycl::minimum<>(),
+  test<class KernelNamePlusV>(input, sycl::plus<>(), 0);
+  test<class KernelNameMinimumV>(input, sycl::minimum<>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumV>(q, input, output, sycl::maximum<>(),
+  test<class KernelNameMaximumV>(input, sycl::maximum<>(),
                                  std::numeric_limits<int>::lowest());
 
-  test<class KernelNamePlusI>(q, input, output, sycl::plus<int>(), 0);
-  test<class KernelNameMinimumI>(q, input, output, sycl::minimum<int>(),
+  test<class KernelNamePlusI>(input, sycl::plus<int>(), 0);
+  test<class KernelNameMinimumI>(input, sycl::minimum<int>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
+  test<class KernelNameMaximumI>(input, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
-  test<class KernelNameMultipliesI>(q, input_small, output_small,
-                                    sycl::multiplies<int>(), 1);
-  test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
-  test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
-  test<class KernelNameBitAndI>(q, input_small, output_small,
-                                sycl::bit_and<int>(), ~0);
+  test<class KernelNameMultipliesI>(input_small, sycl::multiplies<int>(), 1);
+  test<class KernelNameBitOrI>(input, sycl::bit_or<int>(), 0);
+  test<class KernelNameBitXorI>(input, sycl::bit_xor<int>(), 0);
+  test<class KernelNameBitAndI>(input_small, sycl::bit_and<int>(), ~0);
 
-  test<class LogicalOrInt>(q, input, output, sycl::logical_or<int>(), 0);
-  test<class LogicalAndInt>(q, input, output, sycl::logical_and<int>(), 1);
+  test<class LogicalOrInt>(input, sycl::logical_or<int>(), 0);
+  test<class LogicalAndInt>(input, sycl::logical_and<int>(), 1);
 
   std::array<bool, N> bool_input = {};
-  std::array<bool, N> bool_output = {};
-  test<class LogicalOrBool>(q, bool_input, bool_output,
-                            sycl::logical_or<bool>(), false);
-  test<class LogicalOrVoid>(q, bool_input, bool_output, sycl::logical_or<>(),
-                            false);
-  test<class LogicalAndBool>(q, bool_input, bool_output,
-                             sycl::logical_and<bool>(), true);
-  test<class LogicalAndVoid>(q, bool_input, bool_output, sycl::logical_and<>(),
-                             true);
+  test<class LogicalOrBool>(bool_input, sycl::logical_or<bool>(), false);
+  test<class LogicalOrVoid>(bool_input, sycl::logical_or<>(), false);
+  test<class LogicalAndBool>(bool_input, sycl::logical_and<bool>(), true);
+  test<class LogicalAndVoid>(bool_input, sycl::logical_and<>(), true);
+
+  test<class LogicalOrBool>(bool_input, sycl::logical_or<bool>(), false);
+  test<class LogicalOrVoid>(bool_input, sycl::logical_or<>(), false);
+  test<class LogicalAndBool>(bool_input, sycl::logical_and<bool>(), true);
+  test<class LogicalAndVoid>(bool_input, sycl::logical_and<>(), true);
+
+  std::array<int2, N> int2_input = {};
+  std::iota(int2_input.begin(), int2_input.end(), 0);
+  test<class PlusInt2>(int2_input, sycl::plus<int2>(), {0, 0});
+  test<class PlusInt2V>(int2_input, sycl::plus<>(), {0, 0});
+
+  if (q.get_device().has(aspect::fp16)) {
+    std::array<half, 32> half_input = {};
+    std::iota(half_input.begin(), half_input.end(), 0);
+    test<class PlusHalf>(half_input, sycl::plus<half>(), 0);
+    test<class PlusHalfV>(half_input, sycl::plus<>(), 0);
+  }
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_inclusive_scan and
@@ -206,22 +214,17 @@ int main() {
   // sycl::plus binary operation.
 #ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
   std::array<std::complex<float>, N> input_cf;
-  std::array<std::complex<float>, N> output_cf;
   std::iota(input_cf.begin(), input_cf.end(), 0);
-  std::fill(output_cf.begin(), output_cf.end(), 0);
-  test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
+  test<class KernelNamePlusComplexF>(input_cf,
                                      sycl::plus<std::complex<float>>(), 0);
-  test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
+  test<class KernelNamePlusUnspecF>(input_cf, sycl::plus<>(), 0);
 
   if (q.get_device().has(aspect::fp64)) {
     std::array<std::complex<double>, N> input_cd;
-    std::array<std::complex<double>, N> output_cd;
     std::iota(input_cd.begin(), input_cd.end(), 0);
-    std::fill(output_cd.begin(), output_cd.end(), 0);
-    test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
+    test<class KernelNamePlusComplexD>(input_cd,
                                        sycl::plus<std::complex<double>>(), 0);
-    test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
-                                      0);
+    test<class KernelNamePlusUnspecD>(input_cd, sycl::plus<>(), 0);
   } else {
     std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
               << std::endl;
diff --git a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp
index 8e6c9d75a29b..b89032a098fb 100644
--- a/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp
+++ b/sycl/test-e2e/GroupAlgorithm/reduce_sycl2020.cpp
@@ -10,18 +10,21 @@
 #include <limits>
 #include <numeric>
 #include <sycl/sycl.hpp>
+
 using namespace sycl;
 
+queue q;
+
 template <typename SpecializationKernelName, typename InputContainer,
-          typename OutputContainer, class BinaryOperation>
-void test(queue q, InputContainer input, OutputContainer output,
-          BinaryOperation binary_op,
-          typename OutputContainer::value_type identity) {
+          class BinaryOperation>
+void test(const InputContainer &input, BinaryOperation binary_op,
+          typename InputContainer::value_type identity) {
   typedef typename InputContainer::value_type InputT;
-  typedef typename OutputContainer::value_type OutputT;
-  constexpr OutputT init = 42;
-  constexpr size_t N = input.size();
-  constexpr size_t G = 64;
+  typedef InputT OutputT;
+  std::array<OutputT, 6> output = {};
+  constexpr OutputT init(42);
+  size_t N = input.size();
+  constexpr size_t G = 16;
   {
     buffer<InputT> in_buf(input.data(), input.size());
     buffer<OutputT> out_buf(output.data(), output.size());
@@ -51,18 +54,18 @@ void test(queue q, InputContainer input, OutputContainer output,
   }
   // std::reduce is not implemented yet, so use std::accumulate instead
   // TODO: use std::reduce when it will be supported
-  assert(output[0] == std::accumulate(input.begin(), input.begin() + G,
-                                      identity, binary_op));
-  assert(output[1] ==
-         std::accumulate(input.begin(), input.begin() + G, init, binary_op));
-  assert(output[2] ==
-         std::accumulate(input.begin(), input.end(), identity, binary_op));
-  assert(output[3] ==
-         std::accumulate(input.begin(), input.end(), init, binary_op));
-  assert(output[4] ==
-         std::accumulate(input.begin(), input.end(), identity, binary_op));
-  assert(output[5] ==
-         std::accumulate(input.begin(), input.end(), init, binary_op));
+  assert(equal(output[0], std::accumulate(input.begin(), input.begin() + G,
+                                          identity, binary_op)));
+  assert(equal(output[1], std::accumulate(input.begin(), input.begin() + G,
+                                          init, binary_op)));
+  assert(equal(output[2], std::accumulate(input.begin(), input.end(), identity,
+                                          binary_op)));
+  assert(equal(output[3],
+               std::accumulate(input.begin(), input.end(), init, binary_op)));
+  assert(equal(output[4], std::accumulate(input.begin(), input.end(), identity,
+                                          binary_op)));
+  assert(equal(output[5],
+               std::accumulate(input.begin(), input.end(), init, binary_op)));
 }
 
 int main() {
@@ -74,41 +77,45 @@ int main() {
 
   constexpr int N = 128;
   std::array<int, N> input;
-  std::array<int, 6> output;
   std::iota(input.begin(), input.end(), 0);
-  std::fill(output.begin(), output.end(), 0);
 
-  test<class KernelNamePlusV>(q, input, output, sycl::plus<>(), 0);
-  test<class KernelNameMinimumV>(q, input, output, sycl::minimum<>(),
+  test<class KernelNamePlusV>(input, sycl::plus<>(), 0);
+  test<class KernelNameMinimumV>(input, sycl::minimum<>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumV>(q, input, output, sycl::maximum<>(),
+  test<class KernelNameMaximumV>(input, sycl::maximum<>(),
                                  std::numeric_limits<int>::lowest());
 
-  test<class KernelNamePlusI>(q, input, output, sycl::plus<int>(), 0);
-  test<class KernelNameMinimumI>(q, input, output, sycl::minimum<int>(),
+  test<class KernelNamePlusI>(input, sycl::plus<int>(), 0);
+  test<class KernelNameMinimumI>(input, sycl::minimum<int>(),
                                  std::numeric_limits<int>::max());
-  test<class KernelNameMaximumI>(q, input, output, sycl::maximum<int>(),
+  test<class KernelNameMaximumI>(input, sycl::maximum<int>(),
                                  std::numeric_limits<int>::lowest());
 
-  test<class KernelNameMultipliesI>(q, input, output, sycl::multiplies<int>(),
-                                    1);
-  test<class KernelNameBitOrI>(q, input, output, sycl::bit_or<int>(), 0);
-  test<class KernelNameBitXorI>(q, input, output, sycl::bit_xor<int>(), 0);
-  test<class KernelNameBitAndI>(q, input, output, sycl::bit_and<int>(), ~0);
+  test<class KernelNameMultipliesI>(input, sycl::multiplies<int>(), 1);
+  test<class KernelNameBitOrI>(input, sycl::bit_or<int>(), 0);
+  test<class KernelNameBitXorI>(input, sycl::bit_xor<int>(), 0);
+  test<class KernelNameBitAndI>(input, sycl::bit_and<int>(), ~0);
 
-  test<class LogicalOrInt>(q, input, output, sycl::logical_or<int>(), 0);
-  test<class LogicalAndInt>(q, input, output, sycl::logical_and<int>(), 1);
+  test<class LogicalOrInt>(input, sycl::logical_or<int>(), 0);
+  test<class LogicalAndInt>(input, sycl::logical_and<int>(), 1);
 
   std::array<bool, N> bool_input = {};
-  std::array<bool, 6> bool_output = {};
-  test<class LogicalOrBool>(q, bool_input, bool_output,
-                            sycl::logical_or<bool>(), false);
-  test<class LogicalOrVoid>(q, bool_input, bool_output, sycl::logical_or<>(),
-                            false);
-  test<class LogicalAndBool>(q, bool_input, bool_output,
-                             sycl::logical_and<bool>(), true);
-  test<class LogicalAndVoid>(q, bool_input, bool_output, sycl::logical_and<>(),
-                             true);
+  test<class LogicalOrBool>(bool_input, sycl::logical_or<bool>(), false);
+  test<class LogicalOrVoid>(bool_input, sycl::logical_or<>(), false);
+  test<class LogicalAndBool>(bool_input, sycl::logical_and<bool>(), true);
+  test<class LogicalAndVoid>(bool_input, sycl::logical_and<>(), true);
+
+  std::array<int2, N> int2_input = {};
+  std::iota(int2_input.begin(), int2_input.end(), 0);
+  test<class PlusInt2>(int2_input, sycl::plus<int2>(), {0, 0});
+  test<class PlusInt2V>(int2_input, sycl::plus<>(), {0, 0});
+
+  if (q.get_device().has(aspect::fp16)) {
+    std::array<half, 32> half_input = {};
+    std::iota(half_input.begin(), half_input.end(), 0);
+    test<class PlusHalf>(half_input, sycl::plus<half>(), 0);
+    test<class PlusHalfV>(half_input, sycl::plus<>(), 0);
+  }
 
   // as part of SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS (
   // https://github.com/intel/llvm/pull/5108/ ) joint_reduce and
@@ -116,22 +123,17 @@ int main() {
   // sycl::plus binary operation.
 #ifdef SYCL_EXT_ONEAPI_COMPLEX_ALGORITHMS
   std::array<std::complex<float>, N> input_cf;
-  std::array<std::complex<float>, 6> output_cf;
   std::iota(input_cf.begin(), input_cf.end(), 0);
-  std::fill(output_cf.begin(), output_cf.end(), 0);
-  test<class KernelNamePlusComplexF>(q, input_cf, output_cf,
+  test<class KernelNamePlusComplexF>(input_cf,
                                      sycl::plus<std::complex<float>>(), 0);
-  test<class KernelNamePlusUnspecF>(q, input_cf, output_cf, sycl::plus<>(), 0);
+  test<class KernelNamePlusUnspecF>(input_cf, sycl::plus<>(), 0);
 
   if (q.get_device().has(aspect::fp64)) {
     std::array<std::complex<double>, N> input_cd;
-    std::array<std::complex<double>, 6> output_cd;
     std::iota(input_cd.begin(), input_cd.end(), 0);
-    std::fill(output_cd.begin(), output_cd.end(), 0);
-    test<class KernelNamePlusComplexD>(q, input_cd, output_cd,
+    test<class KernelNamePlusComplexD>(input_cd,
                                        sycl::plus<std::complex<double>>(), 0);
-    test<class KernelNamePlusUnspecD>(q, input_cd, output_cd, sycl::plus<>(),
-                                      0);
+    test<class KernelNamePlusUnspecD>(input_cd, sycl::plus<>(), 0);
   } else {
     std::cout << "aspect::fp64 not supported. skipping std::complex<double>"
               << std::endl;
diff --git a/sycl/test-e2e/GroupAlgorithm/support.h b/sycl/test-e2e/GroupAlgorithm/support.h
index ef3be0138d55..9d1f3ffddb3a 100644
--- a/sycl/test-e2e/GroupAlgorithm/support.h
+++ b/sycl/test-e2e/GroupAlgorithm/support.h
@@ -24,3 +24,22 @@ bool isSupportedDevice(device D) {
 
   return false;
 }
+
+template <typename T, typename S> bool equal(const T &x, const S &y) {
+  // vec equal returns a vector of which components were equal
+  if constexpr (sycl::detail::is_vec<T>::value) {
+    for (int i = 0; i < x.size(); ++i)
+      if (x[i] != y[i])
+        return false;
+    return true;
+  } else
+    return x == y;
+}
+
+template <typename T1, typename T2>
+bool ranges_equal(T1 begin1, T1 end1, T2 begin2) {
+  for (; begin1 != end1; ++begin1, ++begin2)
+    if (!equal(*begin1, *begin2))
+      return false;
+  return true;
+}