Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Reduce with Teams broken for custom initialize #407

Closed
khpierson opened this issue Aug 29, 2016 · 0 comments
Closed

Reduce with Teams broken for custom initialize #407

khpierson opened this issue Aug 29, 2016 · 0 comments
Assignees
Labels
Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Milestone

Comments

@khpierson
Copy link

Below is a patch that fixes the issue.

--- a/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_Parallel.hpp
+++ b/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_Parallel.hpp
@@ -1040,8 +1040,10 @@ public:

     pointer_type const result = (pointer_type) (m_unified_space ? m_unified_space : m_scratch_space) ;

+    value_type init;
+    ValueInit::init( ReducerConditional::select(m_functor , m_reducer) , &init);
     if(Impl::cuda_inter_block_reduction<FunctorType,ValueJoin,WorkTag>
-           (value,ValueJoin(ReducerConditional::select(m_functor , m_reducer)),m_scratch_space,result,m_scratch_flags,blockDim.y)) {
+           (value,init,ValueJoin(ReducerConditional::select(m_functor , m_reducer)),m_scratch_space,result,m_scratch_flags,blockDim.y)) {
       const unsigned id = threadIdx.y*blockDim.x + threadIdx.x;
       if(id==0) {
         Kokkos::Impl::FunctorFinal< ReducerTypeFwd , WorkTag >::final( ReducerConditional::select(m_functor , m_reducer) , (void*) &value );
diff --git a/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp b/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp
index cff183f..c9539a3 100644
--- a/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp
+++ b/TPLs_src/Trilinos/packages/kokkos/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp
@@ -133,6 +133,7 @@ inline void cuda_intra_block_reduction( ValueType& value,
 template< class FunctorType , class JoinOp , class ArgTag = void >
 __device__
 bool cuda_inter_block_reduction( typename FunctorValueTraits< FunctorType , ArgTag >::reference_type  value,
+                                 typename FunctorValueTraits< FunctorType , ArgTag >::reference_type  neutral,
                                  const JoinOp& join,
                                  Cuda::size_type * const m_scratch_space,
                                  typename FunctorValueTraits< FunctorType , ArgTag >::pointer_type const result,
@@ -170,7 +171,7 @@ bool cuda_inter_block_reduction( typename FunctorValueTraits< FunctorType , ArgT
       if(id == 0)
         *m_scratch_flags = 0;
       last_block = true;
-      value = value_type();
+      value = neutral;

       pointer_type const volatile global = (pointer_type) m_scratch_space ;
@crtrott crtrott added the Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos) label Aug 29, 2016
@crtrott crtrott added this to the Summer 2016 milestone Aug 29, 2016
@crtrott crtrott self-assigned this Aug 29, 2016
crtrott added a commit that referenced this issue Aug 29, 2016
@crtrott crtrott closed this as completed Sep 3, 2016
@hcedwar hcedwar added this to InProgress in Custom Reductions May 17, 2017
@crtrott crtrott moved this from InProgress to Closed in Custom Reductions May 18, 2017
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Bug Broken / incorrect code; it could be Kokkos' responsibility, or others’ (e.g., Trilinos)
Projects
No open projects
Development

No branches or pull requests

2 participants