From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Thu, 14 Feb 2019 02:31:17 +0000 (+0000)
Subject: Do not use cudaDeviceSynchronize in release mode
X-Git-Tag: v9.1.0-rc1~341^2
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=c46ca30fe9b73ae1c4ac13010f8377705b6e285e;p=dealii.git

Do not use cudaDeviceSynchronize in release mode
---

diff --git a/include/deal.II/lac/vector_operations_internal.h b/include/deal.II/lac/vector_operations_internal.h
index c78df10c94..5d98e3f23a 100644
--- a/include/deal.II/lac/vector_operations_internal.h
+++ b/include/deal.II/lac/vector_operations_internal.h
@@ -2026,10 +2026,12 @@ namespace internal
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::set<Number>
           <<<n_blocks, block_size>>>(data.values_dev.get(), s, size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2049,10 +2051,12 @@ namespace internal
                                      v_data.values_dev.get(),
                                      size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2072,10 +2076,12 @@ namespace internal
                                      v_data.values_dev.get(),
                                      size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2091,10 +2097,12 @@ namespace internal
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add<Number>
           <<<n_blocks, block_size>>>(data.values_dev.get(), a, size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2115,10 +2123,12 @@ namespace internal
                                      v_data.values_dev.get(),
                                      size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2144,10 +2154,12 @@ namespace internal
                                                     w_data.values_dev.get(),
                                                     size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2166,10 +2178,12 @@ namespace internal
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(
             x, data.values_dev.get(), 1., v_data.values_dev.get(), size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2189,10 +2203,12 @@ namespace internal
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(
             x, data.values_dev.get(), a, v_data.values_dev.get(), size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2220,10 +2236,12 @@ namespace internal
                                                     w_data.values_dev.get(),
                                                     size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2239,10 +2257,12 @@ namespace internal
         ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale<Number>
           <<<n_blocks, block_size>>>(data.values_dev.get(), factor, size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2261,10 +2281,12 @@ namespace internal
                                                     v_data.values_dev.get(),
                                                     size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2285,10 +2307,12 @@ namespace internal
                                                     v_data.values_dev.get(),
                                                     size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static void
@@ -2314,10 +2338,12 @@ namespace internal
                                                     w_data.values_dev.get(),
                                                     size);
 
+#  ifdef DEBUG
         // Check that the kernel was launched correctly
         AssertCuda(cudaGetLastError());
         // Check that there was no problem during the execution of the kernel
         AssertCuda(cudaDeviceSynchronize());
+#  endif
       }
 
       static Number
diff --git a/source/lac/cuda_sparse_matrix.cu b/source/lac/cuda_sparse_matrix.cu
index 6756e9d97d..687e565d26 100644
--- a/source/lac/cuda_sparse_matrix.cu
+++ b/source/lac/cuda_sparse_matrix.cu
@@ -343,10 +343,12 @@ namespace CUDAWrappers
     internal::scale<Number>
       <<<n_blocks, block_size>>>(val_dev.get(), factor, nnz);
 
+#  ifdef DEBUG
     // Check that the kernel was launched correctly
     AssertCuda(cudaGetLastError());
     // Check that there was no problem during the execution of the kernel
     AssertCuda(cudaDeviceSynchronize());
+#  endif
 
     return *this;
   }
@@ -363,10 +365,12 @@ namespace CUDAWrappers
     internal::scale<Number>
       <<<n_blocks, block_size>>>(val_dev.get(), 1. / factor, nnz);
 
+#  ifdef DEBUG
     // Check that the kernel was launched correctly
     AssertCuda(cudaGetLastError());
     // Check that there was no problem during the execution of the kernel
     AssertCuda(cudaDeviceSynchronize());
+#  endif
 
     return *this;
   }
@@ -515,10 +519,13 @@ namespace CUDAWrappers
                                  column_index_dev.get(),
                                  row_ptr_dev.get(),
                                  column_sums.get_values());
+
+#  ifdef DEBUG
     // Check that the kernel was launched correctly
     AssertCuda(cudaGetLastError());
     // Check that there was no problem during the execution of the kernel
     AssertCuda(cudaDeviceSynchronize());
+#  endif
 
     return column_sums.linfty_norm();
   }
@@ -537,10 +544,13 @@ namespace CUDAWrappers
                                  column_index_dev.get(),
                                  row_ptr_dev.get(),
                                  row_sums.get_values());
+
+#  ifdef DEUG
     // Check that the kernel was launched correctly
     AssertCuda(cudaGetLastError());
     // Check that there was no problem during the execution of the kernel
     AssertCuda(cudaDeviceSynchronize());
+#  endif
 
     return row_sums.linfty_norm();
   }
diff --git a/source/lac/cuda_vector.cu b/source/lac/cuda_vector.cu
index 4f12a49740..0787cb679c 100644
--- a/source/lac/cuda_vector.cu
+++ b/source/lac/cuda_vector.cu
@@ -158,10 +158,13 @@ namespace LinearAlgebra
 
           kernel::vector_bin_op<Number, kernel::Binop_Addition>
             <<<n_blocks, block_size>>>(val.get(), tmp, n_elements);
+
+#  ifdef DEBUG
           // Check that the kernel was launched correctly
           AssertCuda(cudaGetLastError());
           // Check that there was no problem during the execution of the kernel
           AssertCuda(cudaDeviceSynchronize());
+#  endif
 
           // Delete the temporary vector
           Utilities::CUDA::free(tmp);
@@ -197,10 +200,12 @@ namespace LinearAlgebra
       kernel::vec_scale<Number>
         <<<n_blocks, block_size>>>(val.get(), factor, n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
 
       return *this;
     }
@@ -217,10 +222,12 @@ namespace LinearAlgebra
       kernel::vec_scale<Number>
         <<<n_blocks, block_size>>>(val.get(), 1. / factor, n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
 
       return *this;
     }
@@ -246,10 +253,12 @@ namespace LinearAlgebra
       kernel::vector_bin_op<Number, kernel::Binop_Addition>
         <<<n_blocks, block_size>>>(val.get(), down_V.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
 
       return *this;
     }
@@ -275,10 +284,12 @@ namespace LinearAlgebra
       kernel::vector_bin_op<Number, kernel::Binop_Subtraction>
         <<<n_blocks, block_size>>>(val.get(), down_V.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
 
       return *this;
     }
@@ -336,10 +347,12 @@ namespace LinearAlgebra
       kernel::vec_add<Number>
         <<<n_blocks, block_size>>>(val.get(), a, n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }
 
 
@@ -364,10 +377,12 @@ namespace LinearAlgebra
       kernel::add_aV<Number><<<dim3(n_blocks, 1), dim3(block_size)>>>(
         val.get(), a, down_V.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }
 
 
@@ -406,10 +421,12 @@ namespace LinearAlgebra
       kernel::add_aVbW<Number><<<dim3(n_blocks, 1), dim3(block_size)>>>(
         val.get(), a, down_V.val.get(), b, down_W.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }
 
 
@@ -437,10 +454,12 @@ namespace LinearAlgebra
       kernel::sadd<Number><<<dim3(n_blocks, 1), dim3(block_size)>>>(
         s, val.get(), a, down_V.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }
 
 
@@ -464,10 +483,12 @@ namespace LinearAlgebra
       kernel::scale<Number><<<dim3(n_blocks, 1), dim3(block_size)>>>(
         val.get(), down_scaling_factors.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }
 
 
@@ -493,10 +514,12 @@ namespace LinearAlgebra
       kernel::equ<Number><<<dim3(n_blocks, 1), dim3(block_size)>>>(
         val.get(), a, down_V.val.get(), n_elements);
 
+#  ifdef DEBUG
       // Check that the kernel was launched correctly
       AssertCuda(cudaGetLastError());
       // Check that there was no problem during the execution of the kernel
       AssertCuda(cudaDeviceSynchronize());
+#  endif
     }