From 77270d51bfb527f460eee6d5921bd61640ed157c Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Sun, 26 Apr 2020 19:11:37 -0400 Subject: [PATCH] Introduce AssertCudaKernel to replace two calls to AssertCuda --- include/deal.II/base/exceptions.h | 31 +++++ include/deal.II/base/undefine_macros.h | 4 + .../lac/affine_constraints.templates.h | 7 +- include/deal.II/lac/precondition.h | 8 +- .../deal.II/lac/vector_operations_internal.h | 112 +++--------------- source/lac/cuda_sparse_matrix.cu | 32 +---- source/lac/cuda_vector.cu | 88 ++------------ 7 files changed, 66 insertions(+), 216 deletions(-) diff --git a/include/deal.II/base/exceptions.h b/include/deal.II/base/exceptions.h index 6444344f11..1426b2caf3 100644 --- a/include/deal.II/base/exceptions.h +++ b/include/deal.II/base/exceptions.h @@ -1769,6 +1769,37 @@ namespace internal } # endif +/** + * An assertion that checks that the kernel was launched and executed + * successfully. + * + * @note This and similar macro names are examples of preprocessor definitions + * in the deal.II library that are not prefixed by a string that likely makes + * them unique to deal.II. As a consequence, it is possible that other + * libraries your code interfaces with define the same name, and the result + * will be name collisions (see + * https://en.wikipedia.org/wiki/Name_collision). One can \#undef + * this macro, as well as all other macros defined by deal.II that are not + * prefixed with either DEAL or deal, by including + * the header deal.II/base/undefine_macros.h after all other + * deal.II headers have been included. + * + * @ingroup Exceptions + * @author Bruno Turcksin, 2020 + */ +# ifdef DEBUG +# define AssertCudaKernel() \ + { \ + cudaError_t local_error_code = cudaPeekAtLastError(); \ + AssertCuda(local_error_code); \ + local_error_code = cudaDeviceSynchronize(); \ + AssertCuda(local_error_code) \ + } +# else +# define AssertCudaKernel() \ + {} +# endif + /** * An assertion that checks that the error code produced by calling a cuSPARSE * routine is equal to CUSPARSE_STATUS_SUCCESS. diff --git a/include/deal.II/base/undefine_macros.h b/include/deal.II/base/undefine_macros.h index 9178ae3313..e3828267ea 100644 --- a/include/deal.II/base/undefine_macros.h +++ b/include/deal.II/base/undefine_macros.h @@ -30,6 +30,10 @@ # undef AssertCuda #endif // #ifdef AssertCuda +#ifdef AssertCudaKernel +# undef AssertCudaKernel +#endif // #ifdef AssertCudaKernel + #ifdef AssertCusolver # undef AssertCusolver #endif // #ifdef AssertCusolver diff --git a/include/deal.II/lac/affine_constraints.templates.h b/include/deal.II/lac/affine_constraints.templates.h index d79ea8864a..1b14df9eb6 100644 --- a/include/deal.II/lac/affine_constraints.templates.h +++ b/include/deal.II/lac/affine_constraints.templates.h @@ -1924,12 +1924,7 @@ namespace internal const int n_blocks = 1 + (n_constraints - 1) / CUDAWrappers::block_size; set_zero_kernel<<>>( constrained_local_dofs_device, n_constraints, vec.get_values()); -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); Utilities::CUDA::free(constrained_local_dofs_device); diff --git a/include/deal.II/lac/precondition.h b/include/deal.II/lac/precondition.h index 295965cd63..a5ff23a69c 100644 --- a/include/deal.II/lac/precondition.h +++ b/include/deal.II/lac/precondition.h @@ -2165,13 +2165,7 @@ namespace internal 1 + (n_local_elements - 1) / CUDAWrappers::block_size; set_initial_guess_kernel<<>>( first_local_range, n_local_elements, vector.get_values()); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); const Number mean_value = vector.mean_value(); vector.add(-mean_value); diff --git a/include/deal.II/lac/vector_operations_internal.h b/include/deal.II/lac/vector_operations_internal.h index e047bb23e9..4ca18f75b3 100644 --- a/include/deal.II/lac/vector_operations_internal.h +++ b/include/deal.II/lac/vector_operations_internal.h @@ -2093,13 +2093,7 @@ namespace internal const int n_blocks = 1 + size / (chunk_size * block_size); ::dealii::LinearAlgebra::CUDAWrappers::kernel::set <<>>(data.values_dev.get(), s, size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2118,13 +2112,7 @@ namespace internal 1., v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2143,13 +2131,7 @@ namespace internal -1., v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2164,13 +2146,7 @@ namespace internal const int n_blocks = 1 + size / (chunk_size * block_size); ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_add <<>>(data.values_dev.get(), a, size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2190,13 +2166,7 @@ namespace internal a, v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2221,13 +2191,7 @@ namespace internal b, w_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2245,13 +2209,7 @@ namespace internal ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd <<>>( x, data.values_dev.get(), 1., v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2270,13 +2228,7 @@ namespace internal ::dealii::LinearAlgebra::CUDAWrappers::kernel::sadd <<>>( x, data.values_dev.get(), a, v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2303,13 +2255,7 @@ namespace internal b, w_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2324,13 +2270,7 @@ namespace internal const int n_blocks = 1 + size / (chunk_size * block_size); ::dealii::LinearAlgebra::CUDAWrappers::kernel::vec_scale <<>>(data.values_dev.get(), factor, size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2348,13 +2288,7 @@ namespace internal <<>>(data.values_dev.get(), v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2374,13 +2308,7 @@ namespace internal a, v_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static void @@ -2405,13 +2333,7 @@ namespace internal b, w_data.values_dev.get(), size); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } static Number @@ -2438,13 +2360,7 @@ namespace internal v_data.values_dev.get(), static_cast( size)); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); // Copy the result back to the host Number result; diff --git a/source/lac/cuda_sparse_matrix.cu b/source/lac/cuda_sparse_matrix.cu index 013159bfd3..7e331e5654 100644 --- a/source/lac/cuda_sparse_matrix.cu +++ b/source/lac/cuda_sparse_matrix.cu @@ -342,13 +342,7 @@ namespace CUDAWrappers const int n_blocks = 1 + (nnz - 1) / block_size; internal::scale <<>>(val_dev.get(), factor, nnz); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -364,13 +358,7 @@ namespace CUDAWrappers const int n_blocks = 1 + (nnz - 1) / block_size; internal::scale <<>>(val_dev.get(), 1. / factor, nnz); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -519,13 +507,7 @@ namespace CUDAWrappers column_index_dev.get(), row_ptr_dev.get(), column_sums.get_values()); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return column_sums.linfty_norm(); } @@ -544,13 +526,7 @@ namespace CUDAWrappers column_index_dev.get(), row_ptr_dev.get(), row_sums.get_values()); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return row_sums.linfty_norm(); } diff --git a/source/lac/cuda_vector.cu b/source/lac/cuda_vector.cu index f2f8bc05d0..fe0234e33b 100644 --- a/source/lac/cuda_vector.cu +++ b/source/lac/cuda_vector.cu @@ -158,13 +158,7 @@ namespace LinearAlgebra kernel::vector_bin_op <<>>(val.get(), tmp, n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); // Delete the temporary vector Utilities::CUDA::free(tmp); @@ -199,13 +193,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::vec_scale <<>>(val.get(), factor, n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -221,13 +209,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::vec_scale <<>>(val.get(), 1. / factor, n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -252,13 +234,7 @@ namespace LinearAlgebra kernel::vector_bin_op <<>>(val.get(), down_V.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -283,13 +259,7 @@ namespace LinearAlgebra kernel::vector_bin_op <<>>(val.get(), down_V.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); return *this; } @@ -346,13 +316,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::vec_add <<>>(val.get(), a, n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } @@ -376,13 +340,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::add_aV<<>>( val.get(), a, down_V.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } @@ -420,13 +378,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); 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(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } @@ -453,13 +405,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::sadd<<>>( s, val.get(), a, down_V.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } @@ -482,13 +428,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::scale<<>>( val.get(), down_scaling_factors.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } @@ -513,13 +453,7 @@ namespace LinearAlgebra const int n_blocks = 1 + (n_elements - 1) / (chunk_size * block_size); kernel::equ<<>>( val.get(), a, down_V.val.get(), n_elements); - -# ifdef DEBUG - // Check that the kernel was launched correctly - AssertCuda(cudaPeekAtLastError()); - // Check that there was no problem during the execution of the kernel - AssertCuda(cudaDeviceSynchronize()); -# endif + AssertCudaKernel(); } -- 2.39.5