From 6a2d9f2c73f0ca47aa9bc0a21be63f972016a53f Mon Sep 17 00:00:00 2001
From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Mon, 11 Nov 2019 14:10:07 +0000
Subject: [PATCH] Make variable volatile to avoid incorrect optimization

See https://devtalk.nvidia.com/default/topic/965588/race-condition-within-warp/
---
 include/deal.II/lac/cuda_kernels.templates.h | 10 +++++-----
 1 file changed, 5 insertions(+), 5 deletions(-)

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)
           {
-- 
2.39.5