From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Mon, 11 Nov 2019 14:10:07 +0000 (+0000)
Subject: Make variable volatile to avoid incorrect optimization
X-Git-Tag: v9.2.0-rc1~885^2
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=refs%2Fpull%2F9038%2Fhead;p=dealii.git

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 8c962b0e93..6eb0469f39 100644
--- 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)
           {