[SYCL] Pull oneDPL tuple to use in reduction implementation (#3481)

Using sycl::detail::tuple is a temporary work-around for various
problems caused by using std::tuple:
a) reduction using std::tuple cannot be compiled on Windows because
   std::tuple cannot be copied to DEVICE.
b) internal error in level_zero RT.

The new sycl::detail::tuple class is a very limited version
of oneDPL's implementation of tuple. It includes such functionality:
- convert from std::tuple and to std::tuple
- tie(), get<I>(), tuple_element, make_tuple

This change enables parallel_for() with number of reductions
more than 1 for level_zero and for Windows.

The corresponding changes in LIT tests: intel/llvm-test-suite#194

Signed-off-by: Vyacheslav N Klochkov <[email protected]>

Author: v-klochkov

sycl/include/CL/sycl/ONEAPI/reduction.hpp

@@ -12,6 +12,7 @@
 #include <CL/sycl/ONEAPI/group_algorithm.hpp>
 #include <CL/sycl/accessor.hpp>
 #include <CL/sycl/atomic.hpp>
+#include <CL/sycl/detail/tuple.hpp>
 #include <CL/sycl/handler.hpp>
 #include <CL/sycl/kernel.hpp>
 
@@ -30,6 +31,15 @@ using cl::sycl::detail::is_sgeninteger;
 using cl::sycl::detail::queue_impl;
 using cl::sycl::detail::remove_AS;
 
+// std::tuple seems to be a) too heavy and b) not copyable to device now
+// Thus sycl::detail::tuple is used instead.
+// Switching from sycl::device::tuple to std::tuple can be done by re-defining
+// the ReduTupleT type and makeReduTupleT() function below.
+template <typename... Ts> using ReduTupleT = sycl::detail::tuple<Ts...>;
+template <typename... Ts> ReduTupleT<Ts...> makeReduTupleT(Ts... Elements) {
+  return sycl::detail::make_tuple(Elements...);
+}
+
 __SYCL_EXPORT size_t reduGetMaxWGSize(shared_ptr_class<queue_impl> Queue,
                                       size_t LocalMemBytesPerWorkItem);
 __SYCL_EXPORT size_t reduComputeWGSize(size_t NWorkItems, size_t MaxWGSize,
@@ -1290,7 +1300,7 @@ reduSaveFinalResultToUserMem(handler &CGH, Reduction &Redu) {
 template <typename... Reductions, size_t... Is>
 auto createReduLocalAccs(size_t Size, handler &CGH,
                          std::index_sequence<Is...>) {
-  return std::make_tuple(
+  return makeReduTupleT(
       std::tuple_element_t<Is, std::tuple<Reductions...>>::getReadWriteLocalAcc(
           Size, CGH)...);
 }
@@ -1302,7 +1312,7 @@ template <bool IsOneWG, typename... Reductions, size_t... Is>
 auto createReduOutAccs(size_t NWorkGroups, handler &CGH,
                        std::tuple<Reductions...> &ReduTuple,
                        std::index_sequence<Is...>) {
-  return std::make_tuple(
+  return makeReduTupleT(
       std::get<Is>(ReduTuple).template getWriteMemForPartialReds<IsOneWG>(
           NWorkGroups, CGH)...);
 }
@@ -1314,19 +1324,19 @@ template <typename... Reductions, size_t... Is>
 auto getReadAccsToPreviousPartialReds(handler &CGH,
                                       std::tuple<Reductions...> &ReduTuple,
                                       std::index_sequence<Is...>) {
-  return std::make_tuple(
+  return makeReduTupleT(
       std::get<Is>(ReduTuple).getReadAccToPreviousPartialReds(CGH)...);
 }
 
 template <typename... Reductions, size_t... Is>
-std::tuple<typename Reductions::result_type...>
+ReduTupleT<typename Reductions::result_type...>
 getReduIdentities(std::tuple<Reductions...> &ReduTuple,
                   std::index_sequence<Is...>) {
   return {std::get<Is>(ReduTuple).getIdentity()...};
 }
 
 template <typename... Reductions, size_t... Is>
-std::tuple<typename Reductions::binary_operation...>
+ReduTupleT<typename Reductions::binary_operation...>
 getReduBOPs(std::tuple<Reductions...> &ReduTuple, std::index_sequence<Is...>) {
   return {std::get<Is>(ReduTuple).getBinaryOperation()...};
 }
@@ -1340,8 +1350,8 @@ getInitToIdentityProperties(std::tuple<Reductions...> &ReduTuple,
 
 template <typename... Reductions, size_t... Is>
 std::tuple<typename Reductions::reducer_type...>
-createReducers(std::tuple<typename Reductions::result_type...> Identities,
-               std::tuple<typename Reductions::binary_operation...> BOPsTuple,
+createReducers(ReduTupleT<typename Reductions::result_type...> Identities,
+               ReduTupleT<typename Reductions::binary_operation...> BOPsTuple,
                std::index_sequence<Is...>) {
   return {typename Reductions::reducer_type{std::get<Is>(Identities),
                                             std::get<Is>(BOPsTuple)}...};
@@ -1357,9 +1367,9 @@ void callReduUserKernelFunc(KernelType KernelFunc, nd_item<Dims> NDIt,
 template <bool Pow2WG, typename... LocalAccT, typename... ReducerT,
           typename... ResultT, size_t... Is>
 void initReduLocalAccs(size_t LID, size_t WGSize,
-                       std::tuple<LocalAccT...> LocalAccs,
+                       ReduTupleT<LocalAccT...> LocalAccs,
                        const std::tuple<ReducerT...> &Reducers,
-                       const std::tuple<ResultT...> Identities,
+                       ReduTupleT<ResultT...> Identities,
                        std::index_sequence<Is...>) {
   std::tie(std::get<Is>(LocalAccs)[LID]...) =
       std::make_tuple(std::get<Is>(Reducers).MValue...);
@@ -1375,9 +1385,9 @@ void initReduLocalAccs(size_t LID, size_t WGSize,
 template <bool UniformPow2WG, typename... LocalAccT, typename... InputAccT,
           typename... ResultT, size_t... Is>
 void initReduLocalAccs(size_t LID, size_t GID, size_t NWorkItems, size_t WGSize,
-                       std::tuple<InputAccT...> LocalAccs,
-                       std::tuple<LocalAccT...> InputAccs,
-                       const std::tuple<ResultT...> Identities,
+                       ReduTupleT<InputAccT...> LocalAccs,
+                       ReduTupleT<LocalAccT...> InputAccs,
+                       ReduTupleT<ResultT...> Identities,
                        std::index_sequence<Is...>) {
   // Normally, the local accessors are initialized with elements from the input
   // accessors. The exception is the case when (GID >= NWorkItems), which
@@ -1402,8 +1412,8 @@ void initReduLocalAccs(size_t LID, size_t GID, size_t NWorkItems, size_t WGSize,
 
 template <typename... LocalAccT, typename... BOPsT, size_t... Is>
 void reduceReduLocalAccs(size_t IndexA, size_t IndexB,
-                         std::tuple<LocalAccT...> LocalAccs,
-                         std::tuple<BOPsT...> BOPs,
+                         ReduTupleT<LocalAccT...> LocalAccs,
+                         ReduTupleT<BOPsT...> BOPs,
                          std::index_sequence<Is...>) {
   std::tie(std::get<Is>(LocalAccs)[IndexA]...) =
       std::make_tuple((std::get<Is>(BOPs)(std::get<Is>(LocalAccs)[IndexA],
@@ -1415,8 +1425,8 @@ template <bool Pow2WG, bool IsOneWG, typename... Reductions,
           typename... Ts, size_t... Is>
 void writeReduSumsToOutAccs(
     size_t OutAccIndex, size_t WGSize, std::tuple<Reductions...> *,
-    std::tuple<OutAccT...> OutAccs, std::tuple<LocalAccT...> LocalAccs,
-    std::tuple<BOPsT...> BOPs, std::tuple<Ts...> IdentityVals,
+    ReduTupleT<OutAccT...> OutAccs, ReduTupleT<LocalAccT...> LocalAccs,
+    ReduTupleT<BOPsT...> BOPs, ReduTupleT<Ts...> IdentityVals,
     std::array<bool, sizeof...(Reductions)> IsInitializeToIdentity,
     std::index_sequence<Is...>) {
   // Add the initial value of user's variable to the final result.
@@ -1528,9 +1538,9 @@ void reduCGFuncImpl(handler &CGH, KernelType KernelFunc,
   auto OutAccsTuple =
       createReduOutAccs<IsOneWG>(NWorkGroups, CGH, ReduTuple, ReduIndices);
   auto IdentitiesTuple = getReduIdentities(ReduTuple, ReduIndices);
+  auto BOPsTuple = getReduBOPs(ReduTuple, ReduIndices);
   auto InitToIdentityProps =
       getInitToIdentityProperties(ReduTuple, ReduIndices);
-  auto BOPsTuple = getReduBOPs(ReduTuple, ReduIndices);
 
   using Name = typename get_reduction_main_kernel_name_t<
       KernelName, KernelType, Pow2WG, IsOneWG, decltype(OutAccsTuple)>::name;

sycl/include/CL/sycl/detail/tuple.hpp

@@ -0,0 +1,208 @@
+//== tuple.hpp - limited trivially copy constructible implementation- C++ --==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+// ===--------------------------------------------------------------------=== //
+
+#pragma once
+
+#include <CL/sycl/detail/defines_elementary.hpp>
+
+#include <cassert>
+#include <iterator>
+#include <tuple>
+#include <type_traits>
+
+__SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+template <typename... T> struct tuple;
+
+template <typename T, typename... Ts, std::size_t... Is>
+std::tuple<Ts...> get_tuple_tail_impl(const std::tuple<T, Ts...> &Tuple,
+                                      const std::index_sequence<Is...> &) {
+  return std::tuple<Ts...>(std::get<Is + 1>(Tuple)...);
+}
+
+template <typename T, typename... Ts>
+std::tuple<Ts...> get_tuple_tail(const std::tuple<T, Ts...> &Tuple) {
+  return get_tuple_tail_impl(Tuple, std::make_index_sequence<sizeof...(Ts)>());
+}
+
+template <typename... Ts> constexpr tuple<Ts...> make_tuple(Ts... Args) {
+  return sycl::detail::tuple<Ts...>{Args...};
+}
+
+template <typename... Ts> auto tie(Ts &... Args) {
+  return sycl::detail::tuple<Ts &...>(Args...);
+}
+
+template <std::size_t N, typename T> struct tuple_element;
+
+template <std::size_t N, typename T, typename... Rest>
+struct tuple_element<N, tuple<T, Rest...>>
+    : tuple_element<N - 1, tuple<Rest...>> {};
+
+template <typename T, typename... Rest>
+struct tuple_element<0, tuple<T, Rest...>> {
+  using type = T;
+};
+
+template <std::size_t I, class T>
+using tuple_element_t = typename tuple_element<I, T>::type;
+
+// Functor returning reference to the selected element of the tuple.
+template <size_t N> struct get {
+  template <typename... Ts>
+  constexpr auto operator()(tuple<Ts...> &Tuple) const
+      -> decltype(get<N - 1>()(Tuple.next)) {
+    return get<N - 1>()(Tuple.next);
+  }
+
+  template <typename... Ts>
+  constexpr auto operator()(const tuple<Ts...> &Tuple) const
+      -> decltype(get<N - 1>()(Tuple.next)) {
+    return get<N - 1>()(Tuple.next);
+  }
+};
+
+// Functor returning reference to selected element of the tuple.
+// Specialization stopping the recursion.
+template <> struct get<0> {
+  template <typename... Ts>
+  using ret_type = typename tuple_element<0, tuple<Ts...>>::type;
+
+  template <typename... Ts>
+  constexpr ret_type<Ts...> &operator()(tuple<Ts...> &Tuple) const noexcept {
+    return Tuple.holder.value;
+  }
+
+  template <typename... Ts>
+  constexpr ret_type<Ts...> const &operator()(const tuple<Ts...> &Tuple) const
+      noexcept {
+    return Tuple.holder.value;
+  }
+};
+
+template <typename T> struct TupleValueHolder {
+  TupleValueHolder() = default;
+  TupleValueHolder(const T &Value) : value(Value) {}
+  T value;
+};
+
+// Tuple needs to be trivially_copy_assignable. Define operator= if necessary.
+template <typename T,
+          bool = std::is_trivially_copy_assignable<TupleValueHolder<T>>::value>
+struct TupleCopyAssignableValueHolder : TupleValueHolder<T> {
+  using TupleValueHolder<T>::TupleValueHolder;
+};
+
+template <typename T>
+struct TupleCopyAssignableValueHolder<T, false> : TupleValueHolder<T> {
+  using TupleValueHolder<T>::TupleValueHolder;
+
+  TupleCopyAssignableValueHolder &
+  operator=(const TupleCopyAssignableValueHolder &RHS) {
+    this->value = RHS.value;
+    return *this;
+  }
+};
+
+template <typename T, typename... Ts> struct tuple<T, Ts...> {
+  TupleCopyAssignableValueHolder<T> holder;
+  tuple<Ts...> next;
+
+  using tuple_type = std::tuple<T, Ts...>;
+
+  tuple() = default;
+  tuple(const tuple &) = default;
+  template <typename UT, typename... UTs>
+  tuple(const tuple<UT, UTs...> &RHS)
+      : holder(RHS.holder.value), next(RHS.next) {}
+
+  tuple(const T &Value, const Ts &... Next) : holder(Value), next(Next...) {}
+
+  // required to convert std::tuple to inner tuple in user-provided functor
+  tuple(const std::tuple<T, Ts...> &RHS)
+      : holder(std::get<0>(RHS)), next(get_tuple_tail(RHS)) {}
+
+  // Convert to std::tuple with the same template arguments.
+  operator std::tuple<T, Ts...>() const {
+    return to_std_tuple(*this, std::make_index_sequence<sizeof...(Ts) + 1>());
+  }
+
+  // Convert to std::tuple with different template arguments.
+  template <typename UT, typename... UTs>
+  operator std::tuple<UT, UTs...>() const {
+    return to_std_tuple(static_cast<tuple<UT, UTs...>>(*this),
+                        std::make_index_sequence<sizeof...(Ts) + 1>());
+  }
+
+  template <typename UT, typename... UTs>
+  tuple &operator=(const detail::tuple<UT, UTs...> &RHS) {
+    holder.value = RHS.holder.value;
+    next = RHS.next;
+    return *this;
+  }
+
+  // if T is deduced with reference, compiler generates deleted operator= and,
+  // since "template operator=" is not considered as operator= overload
+  // the deleted operator= has a preference during lookup
+  tuple &operator=(const detail::tuple<T, Ts...> &) = default;
+
+  // Convert std::tuple to sycl::detail::tuple
+  template <typename UT, typename... UTs>
+  tuple &operator=(const std::tuple<UT, UTs...> &RHS) {
+    holder.value = std::get<0>(RHS);
+    next = get_tuple_tail(RHS);
+    return *this;
+  }
+
+  friend bool operator==(const tuple &LHS, const tuple &RHS) {
+    return LHS.holder.value == RHS.holder.value && LHS.next == RHS.next;
+  }
+  friend bool operator!=(const tuple &LHS, const tuple &RHS) {
+    return !(LHS == RHS);
+  }
+
+  template <typename UT, typename... UTs, std::size_t... Is>
+  static std::tuple<UT, UTs...> to_std_tuple(const tuple<UT, UTs...> &Tuple,
+                                             std::index_sequence<Is...>) {
+    return std::tuple<UT, UTs...>(get<Is>()(Tuple)...);
+  }
+};
+
+template <> struct tuple<> {
+  using tuple_type = std::tuple<>;
+
+  tuple() = default;
+  tuple(const tuple &) = default;
+  tuple(const std::tuple<> &) {}
+
+  tuple &operator=(const tuple &) = default;
+  tuple &operator=(const std::tuple<> &) { return *this; }
+  friend bool operator==(const tuple &, const tuple &) { return true; }
+};
+
+} // namespace detail
+} // namespace sycl
+} // __SYCL_INLINE_NAMESPACE(cl)
+
+namespace std {
+
+template <size_t I, typename... Types>
+constexpr typename tuple_element<I, tuple<Types...>>::type &
+get(cl::sycl::detail::tuple<Types...> &Arg) noexcept {
+  return cl::sycl::detail::get<I>()(Arg);
+}
+
+template <size_t I, typename... Types>
+constexpr typename tuple_element<I, tuple<Types...>>::type const &
+get(const cl::sycl::detail::tuple<Types...> &Arg) noexcept {
+  return cl::sycl::detail::get<I>()(Arg);
+}
+
+} // namespace std