From 6a2d9f2c73f0ca47aa9bc0a21be63f972016a53f Mon Sep 17 00:00:00 2001 From: Bruno Turcksin 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 __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