Address issue kokkos#2989.

Also probably fixes kokkos/kokkos-kernels#645 and kokkos/kokkos-kernels#704. The reproducer in kokkos#2989 now passes with no errors.
crtrott · May 5, 2020 · 2be028b · 2be028b
1 parent 64e9b86
commit 2be028b
Showing 1 changed file with 6 additions and 2 deletions.
diff --git a/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp b/core/src/Cuda/Kokkos_Cuda_ReduceScan.hpp
@@ -97,7 +97,9 @@ __device__ inline
   // Depending on the ValueType _shared__ memory must be aligned up to 8byte
   // boundaries The reason not to use ValueType directly is that for types with
   // constructors it could lead to race conditions
-  __shared__ double sh_result[(sizeof(ValueType) + 7) / 8 * STEP_WIDTH];
+  alignas(alignof(ValueType) > alignof(double) ? alignof(ValueType)
+                                               : alignof(double))
+      __shared__ double sh_result[(sizeof(ValueType) + 7) / 8 * STEP_WIDTH];
   ValueType* result = (ValueType*)&sh_result;
   const int step    = 32 / blockDim.x;
   int shift         = STEP_WIDTH;
@@ -282,7 +284,9 @@ __device__ inline
   // Depending on the ValueType _shared__ memory must be aligned up to 8byte
   // boundaries The reason not to use ValueType directly is that for types with
   // constructors it could lead to race conditions
-  __shared__ double sh_result[(sizeof(ValueType) + 7) / 8 * STEP_WIDTH];
+  alignas(alignof(ValueType) > alignof(double) ? alignof(ValueType)
+                                               : alignof(double))
+      __shared__ double sh_result[(sizeof(ValueType) + 7) / 8 * STEP_WIDTH];
   ValueType* result = (ValueType*)&sh_result;
   const int step    = 32 / blockDim.x;
   int shift         = STEP_WIDTH;