Fix bogus warnings in nested CUDA parallel_reduce

Co-authored-by: Damien L-G <dalg24@gmail.com>
nliber · May 24, 2023 · d92158c · d92158c
1 parent 87c7be9
commit d92158c
Show file tree

Hide file tree

Showing 2 changed files with 19 additions and 13 deletions.
diff --git a/core/src/Cuda/Kokkos_Cuda_Team.hpp b/core/src/Cuda/Kokkos_Cuda_Team.hpp
@@ -486,9 +486,6 @@ KOKKOS_INLINE_FUNCTION std::enable_if_t<Kokkos::is_reducer<ReducerType>::value>
 parallel_reduce(const Impl::TeamThreadRangeBoundariesStruct<
                     iType, Impl::CudaTeamMember>& loop_boundaries,
                 const Closure& closure, const ReducerType& reducer) {
-  (void)loop_boundaries;
-  (void)closure;
-  (void)reducer;
   KOKKOS_IF_ON_DEVICE(
       (typename ReducerType::value_type value;
 
@@ -498,6 +495,11 @@ parallel_reduce(const Impl::TeamThreadRangeBoundariesStruct<
             i < loop_boundaries.end; i += blockDim.y) { closure(i, value); }
 
        loop_boundaries.member.team_reduce(reducer, value);))
+  // Avoid bogus warning about reducer value being uninitialized with combined
+  // reducers
+  KOKKOS_IF_ON_HOST(((void)loop_boundaries; (void)closure;
+                     reducer.init(reducer.reference());
+                     Kokkos::abort("Should only run on the device!");));
 }
 
 /** \brief  Inter-thread parallel_reduce assuming summation.
@@ -546,9 +548,6 @@ KOKKOS_INLINE_FUNCTION std::enable_if_t<Kokkos::is_reducer<ReducerType>::value>
 parallel_reduce(const Impl::TeamVectorRangeBoundariesStruct<
                     iType, Impl::CudaTeamMember>& loop_boundaries,
                 const Closure& closure, const ReducerType& reducer) {
-  (void)loop_boundaries;
-  (void)closure;
-  (void)reducer;
   KOKKOS_IF_ON_DEVICE((typename ReducerType::value_type value;
                        reducer.init(value);
 
@@ -559,6 +558,11 @@ parallel_reduce(const Impl::TeamVectorRangeBoundariesStruct<
 
                        loop_boundaries.member.vector_reduce(reducer, value);
                        loop_boundaries.member.team_reduce(reducer, value);))
+  // Avoid bogus warning about reducer value being uninitialized with combined
+  // reducers
+  KOKKOS_IF_ON_HOST(((void)loop_boundaries; (void)closure;
+                     reducer.init(reducer.reference());
+                     Kokkos::abort("Should only run on the device!");));
 }
 
 template <typename iType, class Closure, typename ValueType>
@@ -626,9 +630,6 @@ KOKKOS_INLINE_FUNCTION std::enable_if_t<is_reducer<ReducerType>::value>
 parallel_reduce(Impl::ThreadVectorRangeBoundariesStruct<
                     iType, Impl::CudaTeamMember> const& loop_boundaries,
                 Closure const& closure, ReducerType const& reducer) {
-  (void)loop_boundaries;
-  (void)closure;
-  (void)reducer;
   KOKKOS_IF_ON_DEVICE((
 
       reducer.init(reducer.reference());
@@ -640,6 +641,11 @@ parallel_reduce(Impl::ThreadVectorRangeBoundariesStruct<
       Impl::CudaTeamMember::vector_reduce(reducer);
 
       ))
+  // Avoid bogus warning about reducer value being uninitialized with combined
+  // reducers
+  KOKKOS_IF_ON_HOST(((void)loop_boundaries; (void)closure;
+                     reducer.init(reducer.reference());
+                     Kokkos::abort("Should only run on the device!");));
 }
 
 /** \brief  Intra-thread vector parallel_reduce.

diff --git a/core/src/impl/Kokkos_Combined_Reducer.hpp b/core/src/impl/Kokkos_Combined_Reducer.hpp
@@ -505,8 +505,8 @@ KOKKOS_INLINE_FUNCTION void parallel_reduce_combined_reducers_impl(
     ReturnTypes&&... returnTypes) noexcept {
   using mem_space_type = typename MemberType::execution_space::memory_space;
 
-  auto combined_value = Impl::make_combined_reducer_value<mem_space_type>(
-      returnType1, returnType2, returnTypes...);
+  decltype(Impl::make_combined_reducer_value<mem_space_type>(
+      returnType1, returnType2, returnTypes...)) combined_value;
 
   auto combined_functor = Impl::make_wrapped_combined_functor<mem_space_type>(
       functor, returnType1, returnType2, returnTypes...);
@@ -543,8 +543,8 @@ auto parallel_reduce(std::string const& label, PolicyType const& policy,
   // directly
   using space_type = Kokkos::DefaultHostExecutionSpace::memory_space;
 
-  auto value = Impl::make_combined_reducer_value<space_type>(
-      returnType1, returnType2, returnTypes...);
+  decltype(Impl::make_combined_reducer_value<space_type>(
+      returnType1, returnType2, returnTypes...)) value;
 
   using combined_reducer_type = Impl::CombinedReducer<
       space_type, Impl::_reducer_from_arg_t<space_type, ReturnType1>,