From c46ca30fe9b73ae1c4ac13010f8377705b6e285e Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Thu, 14 Feb 2019 02:31:17 +0000 Subject: [PATCH] Do not use cudaDeviceSynchronize in release mode --- .../deal.II/lac/vector_operations_internal.h | 26 +++++++++++++++++++ source/lac/cuda_sparse_matrix.cu | 10 +++++++ source/lac/cuda_vector.cu | 23 ++++++++++++++++ 3 files changed, 59 insertions(+) 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 <<>>(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 <<>>(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 <<>>( 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 <<>>( 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 <<>>(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 <<>>(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 <<>>(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 <<>>(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 <<>>(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 <<>>(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 <<>>(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 <<>>(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 <<>>(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<<>>( 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<<>>( 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<<>>( 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<<>>( 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<<>>( 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 } -- 2.39.5