Fixed errors from CI testing

ldh4 · Apr 9, 2024 · 05b6edf · 05b6edf
1 parent cd00bd8
commit 05b6edf
Show file tree

Hide file tree

Showing 3 changed files with 22 additions and 23 deletions.
diff --git a/core/src/Cuda/Kokkos_Cuda_Team.hpp b/core/src/Cuda/Kokkos_Cuda_Team.hpp
@@ -533,11 +533,10 @@ parallel_reduce(const Impl::TeamThreadRangeBoundariesStruct<
       functor_analysis_type::has_join_member_function &&
       functor_analysis_type::has_init_member_function;
 
-  using ReducerSelector =
-      typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
-                                  Sum<ValueType>>::type;
-
   KOKKOS_IF_ON_DEVICE((
+      using ReducerSelector =
+          typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
+                                      Sum<ValueType>>::type;
       auto run_closure =
           [&](ValueType& value) {
             for (iType i = loop_boundaries.start + threadIdx.y;
@@ -614,11 +613,10 @@ parallel_reduce(const Impl::TeamVectorRangeBoundariesStruct<
       functor_analysis_type::has_join_member_function &&
       functor_analysis_type::has_init_member_function;
 
-  using ReducerSelector =
-      typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
-                                  Sum<ValueType>>::type;
-
   KOKKOS_IF_ON_DEVICE((
+      using ReducerSelector =
+          typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
+                                      Sum<ValueType>>::type;
       auto run_closure =
           [&](ValueType& value) {
             for (iType i = loop_boundaries.start + threadIdx.y * blockDim.x +
@@ -735,11 +733,10 @@ parallel_reduce(Impl::ThreadVectorRangeBoundariesStruct<
       functor_analysis_type::has_join_member_function &&
       functor_analysis_type::has_init_member_function;
 
-  using ReducerSelector =
-      typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
-                                  Sum<ValueType>>::type;
-
   KOKKOS_IF_ON_DEVICE((
+      using ReducerSelector =
+          typename Kokkos::Impl::if_c<is_reducer_closure, Closure,
+                                      Sum<ValueType>>::type;
       auto run_closure =
           [&](ValueType& value) {
             for (iType i = loop_boundaries.start + threadIdx.x;

diff --git a/core/src/HIP/Kokkos_HIP_Team.hpp b/core/src/HIP/Kokkos_HIP_Team.hpp
@@ -178,17 +178,19 @@ class HIPTeamMember {
   template <typename ReducerType, typename ValueType>
   KOKKOS_INLINE_FUNCTION void team_reduce(ReducerType const& reducer,
                                           ValueType& value) const noexcept {
-    KOKOS_IF_ON_DEVICE((
-        typename Kokkos::Impl::FunctorAnalysis<
-            FunctorPatternInterface::REDUCE, TeamPolicy<HIP>, ReducerType,
-            typename ReducerType::value_type>::Reducer wrapped_reducer(reducer);
-        hip_intra_block_shuffle_reduction(value, wrapped_reducer, blockDim.y);
-
-        if constexpr (is_reducer_v<ReducerType>) {
-          reducer.reference() = value;
-        }))
+#ifdef __HIP_DEVICE_COMPILE__
+    typename Kokkos::Impl::FunctorAnalysis<
+        FunctorPatternInterface::REDUCE, TeamPolicy<HIP>, ReducerType,
+        ValueType>::Reducer wrapped_reducer(reducer);
+    hip_intra_block_shuffle_reduction(value, wrapped_reducer, blockDim.y);
 
-    KOKKOS_IF_ON_HOST(((void)reducer; (void)value;))
+    if constexpr (is_reducer_v<ReducerType>) {
+      reducer.reference() = value;
+    }
+#else
+    (void)reducer;
+    (void)value;
+#endif
   }
 
   //--------------------------------------------------------------------------

diff --git a/core/src/SYCL/Kokkos_SYCL_Team.hpp b/core/src/SYCL/Kokkos_SYCL_Team.hpp
@@ -123,7 +123,7 @@ class SYCLTeamMember {
   template <typename ReducerType, typename ValueType>
   KOKKOS_INLINE_FUNCTION void team_reduce(ReducerType const& reducer,
                                           ValueType& value) const noexcept {
-    using value_type = typename ValueType;
+    using value_type = ValueType;
 
     auto sg                       = m_item.get_sub_group();
     const auto sub_group_range    = sg.get_local_range()[0];