Introduce AssertCudaKernel to replace two calls to AssertCuda

diff --git a/include/deal.II/base/exceptions.h b/include/deal.II/base/exceptions.h
index 6444344f11d1a5fd3eb700388530726b307b939a..1426b2caf3c66a806843c1b83e4a1190896bc830 100644 (file)
--- 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 <code>\#undef</code>
+ * this macro, as well as all other macros defined by deal.II that are not
+ * prefixed with either <code>DEAL</code> or <code>deal</code>, by including
+ * the header <code>deal.II/base/undefine_macros.h</code> 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 9178ae3313d7f77d5e0e6861278f289588d06438..e3828267ea6036b1830e6e15ced054a40c7862fc 100644 (file)
--- 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 d79ea8864ae6eadf400e1e58e9894f455d4e6f9c..1b14df9eb669684023d11de49d9d75e15fb212db 100644 (file)
--- 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<<<n_blocks, CUDAWrappers::block_size>>>(
         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 295965cd63b0e4ba866321939503b38da158cbea..a5ff23a69c155212e73804aa08607a33960c27a4 100644 (file)
--- 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<<<n_blocks, CUDAWrappers::block_size>>>(
         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 e047bb23e98912c4ff8e2e264c1af1e0b7ac9e55..4ca18f75b377d058376e1be61a81e7c125a37ef7 100644 (file)
--- 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<Number>
           <<<n_blocks, block_size>>>(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<Number>
           <<<n_blocks, block_size>>>(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<Number>
           <<<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(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<Number>
           <<<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(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<Number>
           <<<n_blocks, block_size>>>(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
           <<<dim3(n_blocks, 1), dim3(block_size)>>>(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<unsigned int>(
                                                       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 013159bfd3402c0c1e46b927a71ef05020ab6116..7e331e56546f16066d2a38af91e0360e77b600b6 100644 (file)
--- 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<Number>
       <<<n_blocks, block_size>>>(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<Number>
       <<<n_blocks, block_size>>>(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 f2f8bc05d0ab5cc826888e92e7f6ba7dcf274835..fe0234e33b7b83b1b30b42e1d1d09a064c646ee3 100644 (file)
--- a/source/lac/cuda_vector.cu
+++ b/source/lac/cuda_vector.cu
@@ -158,13 +158,7 @@ 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(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<Number>
         <<<n_blocks, block_size>>>(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<Number>
         <<<n_blocks, block_size>>>(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<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(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<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(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<Number>
         <<<n_blocks, block_size>>>(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<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(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<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(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<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(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<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(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<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(cudaPeekAtLastError());
-      // Check that there was no problem during the execution of the kernel
-      AssertCuda(cudaDeviceSynchronize());
-#  endif
+      AssertCudaKernel();
     }