Merge pull request kokkos#1415 from kokkos/issue-1395

Fix Volta failures other than tasking see Issue 1395
ndellingwood · Feb 18, 2018 · 91e8fa0 · 91e8fa0
2 parents 17fba77 + 043ae39
commit 91e8fa0
Show file tree

Hide file tree

Showing 6 changed files with 21 additions and 10 deletions.
diff --git a/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp b/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp
@@ -512,12 +512,18 @@ void cuda_intra_block_reduce_scan( const FunctorType & functor ,
   const pointer_type tdata_intra = base_data + value_count * threadIdx.y ;
 
   { // Intra-warp reduction:
+    KOKKOS_IMPL_CUDA_SYNCWARP;
     BLOCK_REDUCE_STEP(rtid_intra,tdata_intra,0)
+    KOKKOS_IMPL_CUDA_SYNCWARP;
     BLOCK_REDUCE_STEP(rtid_intra,tdata_intra,1)
+    KOKKOS_IMPL_CUDA_SYNCWARP;
     BLOCK_REDUCE_STEP(rtid_intra,tdata_intra,2)
+    KOKKOS_IMPL_CUDA_SYNCWARP;
     BLOCK_REDUCE_STEP(rtid_intra,tdata_intra,3)
+    KOKKOS_IMPL_CUDA_SYNCWARP;
     BLOCK_REDUCE_STEP(rtid_intra,tdata_intra,4)
-  }
+    KOKKOS_IMPL_CUDA_SYNCWARP;
+ }
 
   __syncthreads(); // Wait for all warps to reduce
 

diff --git a/core/src/Cuda/Kokkos_Cuda_Version_9_8_Compatibility.hpp b/core/src/Cuda/Kokkos_Cuda_Version_9_8_Compatibility.hpp
@@ -1,10 +1,12 @@
 #include<Kokkos_Macros.hpp>
 #if ( CUDA_VERSION < 9000 )
+#define KOKKOS_IMPL_CUDA_SYNCWARP (void) __ballot(1)
 #define KOKKOS_IMPL_CUDA_BALLOT(x) __ballot(x)
 #define KOKKOS_IMPL_CUDA_SHFL(x,y,z) __shfl(x,y,z)
 #define KOKKOS_IMPL_CUDA_SHFL_UP(x,y,z) __shfl_up(x,y,z)
 #define KOKKOS_IMPL_CUDA_SHFL_DOWN(x,y,z) __shfl_down(x,y,z)
 #else
+#define KOKKOS_IMPL_CUDA_SYNCWARP __syncwarp()
 #define KOKKOS_IMPL_CUDA_BALLOT(x) __ballot_sync(0xffffffff,x)
 #define KOKKOS_IMPL_CUDA_SHFL(x,y,z) __shfl_sync(0xffffffff,x,y,z)
 #define KOKKOS_IMPL_CUDA_SHFL_UP(x,y,z) __shfl_up_sync(0xffffffff,x,y,z)

diff --git a/core/src/Kokkos_Macros.hpp b/core/src/Kokkos_Macros.hpp
@@ -170,6 +170,7 @@
   #define KOKKOS_INLINE_FUNCTION       __device__  __host__  inline
   #define KOKKOS_FUNCTION              __device__  __host__
   #ifdef KOKKOS_COMPILER_CLANG
+  #define KOKKOS_INLINE_FUNCTION_DEFAULTED KOKKOS_INLINE_FUNCTION
   #define KOKKOS_FUNCTION_DEFAULTED KOKKOS_FUNCTION
   #endif
 #endif // #if defined( __CUDA_ARCH__ )
@@ -180,6 +181,7 @@
   #define KOKKOS_INLINE_FUNCTION       __attribute__((amp,cpu)) inline
   #define KOKKOS_FUNCTION              __attribute__((amp,cpu))
   #define KOKKOS_LAMBDA                [=] __attribute__((amp,cpu))
+  #define KOKKOS_INLINE_FUNCTION_DEFAULTED KOKKOS_INLINE_FUNCTION
   #define KOKKOS_FUNCTION_DEFAULTED    KOKKOS_FUNCTION
 #endif
 
@@ -416,6 +418,7 @@
 #endif
 
 #if !defined( KOKKOS_FUNCTION_DEFAULTED )
+  #define KOKKOS_INLINE_FUNCTION_DEFAULTED inline
   #define KOKKOS_FUNCTION_DEFAULTED /**/
 #endif
 

diff --git a/core/src/Kokkos_MemoryPool.hpp b/core/src/Kokkos_MemoryPool.hpp
@@ -258,10 +258,10 @@ class MemoryPool {
 
   //--------------------------------------------------------------------------
 
-  KOKKOS_INLINE_FUNCTION MemoryPool( MemoryPool && ) = default ;
-  KOKKOS_INLINE_FUNCTION MemoryPool( const MemoryPool & ) = default ;
-  KOKKOS_INLINE_FUNCTION MemoryPool & operator = ( MemoryPool && ) = default ;
-  KOKKOS_INLINE_FUNCTION MemoryPool & operator = ( const MemoryPool & ) = default ;
+  KOKKOS_INLINE_FUNCTION_DEFAULTED MemoryPool( MemoryPool && ) = default ;
+  KOKKOS_INLINE_FUNCTION_DEFAULTED MemoryPool( const MemoryPool & ) = default ;
+  KOKKOS_INLINE_FUNCTION_DEFAULTED MemoryPool & operator = ( MemoryPool && ) = default ;
+  KOKKOS_INLINE_FUNCTION_DEFAULTED MemoryPool & operator = ( const MemoryPool & ) = default ;
 
   KOKKOS_INLINE_FUNCTION MemoryPool()
     : m_tracker()

diff --git a/core/src/Kokkos_TaskScheduler.hpp b/core/src/Kokkos_TaskScheduler.hpp
@@ -418,16 +418,16 @@ class TaskScheduler
   KOKKOS_INLINE_FUNCTION
   TaskScheduler() : m_track(), m_queue(0) {}
 
-  KOKKOS_INLINE_FUNCTION
+  KOKKOS_INLINE_FUNCTION_DEFAULTED
   TaskScheduler( TaskScheduler && rhs ) = default ;
 
-  KOKKOS_INLINE_FUNCTION
+  KOKKOS_INLINE_FUNCTION_DEFAULTED
   TaskScheduler( TaskScheduler const & rhs ) = default ;
 
-  KOKKOS_INLINE_FUNCTION
+  KOKKOS_INLINE_FUNCTION_DEFAULTED
   TaskScheduler & operator = ( TaskScheduler && rhs ) = default ;
 
-  KOKKOS_INLINE_FUNCTION
+  KOKKOS_INLINE_FUNCTION_DEFAULTED
   TaskScheduler & operator = ( TaskScheduler const & rhs ) = default ;
 
   TaskScheduler( memory_pool const & arg_memory_pool )

diff --git a/core/src/impl/Kokkos_TaskQueue.hpp b/core/src/impl/Kokkos_TaskQueue.hpp
@@ -180,7 +180,7 @@ class TaskBase< void , void , void >
   TaskBase & operator = ( TaskBase && ) = delete ;
   TaskBase & operator = ( const TaskBase & ) = delete ;
 
-  KOKKOS_INLINE_FUNCTION ~TaskBase() = default ;
+  KOKKOS_INLINE_FUNCTION_DEFAULTED ~TaskBase() = default ;
 
   KOKKOS_INLINE_FUNCTION constexpr
   TaskBase()