Make variable volatile to avoid incorrect optimization

See https://devtalk.nvidia.com/default/topic/965588/race-condition-within-warp/

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index 8c962b0e93b84f5a715ddf99ecb73659996ec15e..6eb0469f39d37cd4039d8a2dcabfac59e5d97969 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -191,11 +191,11 @@ namespace LinearAlgebra
 
       template <typename Number, typename Operation>
       __device__ void
-      reduce(Number *        result,
-             Number *        result_buffer,
-             const size_type local_idx,
-             const size_type global_idx,
-             const size_type N)
+      reduce(Number *         result,
+             volatile Number *result_buffer,
+             const size_type  local_idx,
+             const size_type  global_idx,
+             const size_type  N)
       {
         for (size_type s = block_size / 2; s > warp_size; s = s >> 1)
           {