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 c78df10c9433e014bff3ae4d6949dd267617ff0f..5d98e3f23a6bec22be62e84d9b448ccc8023d2e8 100644 (file)
--- 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 6756e9d97dedc25a42a0cfb2ba8e5efae5412185..687e565d26ecb1fe1ed4e7a80695b4835860beab 100644 (file)
--- 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 4f12a49740dbe8185066c28798a95ba8253bffad..0787cb679c708d34b4e3a1cdbc56493f9170e63a 100644 (file)
--- 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
     }