From: Bruno Turcksin 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=6a2d9f2c73f0ca47aa9bc0a21be63f972016a53f;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 __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) {