Fix warnings with Clang-15+Cuda

diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h
index 178ad998d7c0727d5c657b31b5b39068f6a4fa1c..fce8598b561e17676b046c634432fa77aeafb8ea 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.h
+++ b/include/deal.II/lac/cuda_kernels.h
@@ -103,7 +103,8 @@ namespace LinearAlgebra
       struct Binop_Subtraction<std::complex<Number>>
       {
         __device__ static inline std::complex<Number>
-        operation(const std::complex<Number> a, const std::complex<Number> b)
+        operation(const std::complex<Number> a,
+                  const std::complex<Number> /*b*/)
         {
           printf("This function is not implemented for std::complex<Number>!");
           assert(false);

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index 86136320b57ee6525865572a893553e9eb359815..f80b15dba5aac48833004c937157514ea4945af8 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -196,8 +196,8 @@ namespace LinearAlgebra
       reduce(Number *         result,
              volatile Number *result_buffer,
              const size_type  local_idx,
-             const size_type  global_idx,
-             const size_type  N)
+             const size_type /*global_idx*/,
+             const size_type /*N*/)
       {
         for (size_type s = block_size / 2; s > warp_size; s = s >> 1)
           {

diff --git a/include/deal.II/lac/cuda_sparse_matrix.h b/include/deal.II/lac/cuda_sparse_matrix.h
index 166c11ffe72fcda9a5a435d25a499f42fb7e35a4..da735bf9ec9e87dea6af0f11eb950ac16b4f587d 100644 (file)
--- a/include/deal.II/lac/cuda_sparse_matrix.h
+++ b/include/deal.II/lac/cuda_sparse_matrix.h
@@ -377,7 +377,7 @@ namespace CUDAWrappers
 
 
   template <typename Number>
-  inline SparseMatrix<Number>::size_type
+  inline typename SparseMatrix<Number>::size_type
   SparseMatrix<Number>::m() const
   {
     return n_rows;
@@ -386,7 +386,7 @@ namespace CUDAWrappers
 
 
   template <typename Number>
-  inline SparseMatrix<Number>::size_type
+  inline typename SparseMatrix<Number>::size_type
   SparseMatrix<Number>::n() const
   {
     return n_cols;

diff --git a/include/deal.II/matrix_free/cuda_fe_evaluation.h b/include/deal.II/matrix_free/cuda_fe_evaluation.h
index ed694816b61e6c1658efe05e6a8a35e80d064140..3fe5af5aeb6a30b7d15b7ff7afff38895ae5177c 100644 (file)
--- a/include/deal.II/matrix_free/cuda_fe_evaluation.h
+++ b/include/deal.II/matrix_free/cuda_fe_evaluation.h
@@ -399,7 +399,8 @@ namespace CUDAWrappers
       }
     else if (integrate_grad == true)
       {
-        evaluator_tensor_product.integrate_gradient<false>(values, gradients);
+        evaluator_tensor_product.template integrate_gradient<false>(values,
+                                                                    gradients);
         __syncthreads();
       }
   }

diff --git a/include/deal.II/matrix_free/cuda_matrix_free.templates.h b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
index bc5efd162e5d5efbd2a9ea14b3270829b7cf0a76..b08f6673ba0e31a0984907dafd3deebf248b0a31 100644 (file)
--- a/include/deal.II/matrix_free/cuda_matrix_free.templates.h
+++ b/include/deal.II/matrix_free/cuda_matrix_free.templates.h
@@ -261,15 +261,15 @@ namespace CUDAWrappers
       const DoFHandler<dim> &dof_handler,
       const UpdateFlags &    update_flags)
       : data(data)
-      , fe_degree(data->fe_degree)
-      , dofs_per_cell(data->dofs_per_cell)
-      , q_points_per_cell(data->q_points_per_cell)
       , fe_values(mapping,
                   fe,
                   Quadrature<dim>(quad),
                   update_inverse_jacobians | update_quadrature_points |
                     update_values | update_gradients | update_JxW_values)
       , lexicographic_inv(shape_info.lexicographic_numbering)
+      , fe_degree(data->fe_degree)
+      , dofs_per_cell(data->dofs_per_cell)
+      , q_points_per_cell(data->q_points_per_cell)
       , update_flags(update_flags)
       , padding_length(data->get_padding_length())
       , hanging_nodes(dof_handler.get_triangulation())
@@ -323,10 +323,10 @@ namespace CUDAWrappers
       // Setup kernel parameters
       const double apply_n_blocks = std::ceil(
         static_cast<double>(n_cells) / static_cast<double>(cells_per_block));
-      const unsigned int apply_x_n_blocks =
-        std::round(std::sqrt(apply_n_blocks));
-      const unsigned int apply_y_n_blocks =
-        std::ceil(apply_n_blocks / static_cast<double>(apply_x_n_blocks));
+      const auto apply_x_n_blocks =
+        static_cast<unsigned int>(std::round(std::sqrt(apply_n_blocks)));
+      const auto apply_y_n_blocks = static_cast<unsigned int>(
+        std::ceil(apply_n_blocks / static_cast<double>(apply_x_n_blocks)));
 
       data->grid_dim[color] = dim3(apply_x_n_blocks, apply_y_n_blocks);
 
@@ -621,10 +621,10 @@ namespace CUDAWrappers
 
   template <int dim, typename Number>
   MatrixFree<dim, Number>::MatrixFree()
-    : n_dofs(0)
+    : my_id(-1)
+    , n_dofs(0)
     , constrained_dofs(nullptr)
     , padding_length(0)
-    , my_id(-1)
     , dof_handler(nullptr)
   {}
 
@@ -708,7 +708,7 @@ namespace CUDAWrappers
 
 
   template <int dim, typename Number>
-  MatrixFree<dim, Number>::Data
+  typename MatrixFree<dim, Number>::Data
   MatrixFree<dim, Number>::get_data(unsigned int color) const
   {
     Data data_copy;
@@ -936,7 +936,7 @@ namespace CUDAWrappers
                        std::ceil(dim * std::log2(fe_degree + 1.)));
 
     dofs_per_cell     = fe.n_dofs_per_cell();
-    q_points_per_cell = std::pow(n_q_points_1d, dim);
+    q_points_per_cell = static_cast<unsigned int>(std::pow(n_q_points_1d, dim));
 
     const ::dealii::internal::MatrixFreeFunctions::ShapeInfo<Number> shape_info(
       quad, fe);
@@ -1106,14 +1106,14 @@ namespace CUDAWrappers
 
     if (n_constrained_dofs != 0)
       {
-        const unsigned int constraint_n_blocks =
+        const auto constraint_n_blocks = static_cast<unsigned int>(
           std::ceil(static_cast<double>(n_constrained_dofs) /
-                    static_cast<double>(block_size));
-        const unsigned int constraint_x_n_blocks =
-          std::round(std::sqrt(constraint_n_blocks));
-        const unsigned int constraint_y_n_blocks =
+                    static_cast<double>(block_size)));
+        const auto constraint_x_n_blocks =
+          static_cast<unsigned int>(std::round(std::sqrt(constraint_n_blocks)));
+        const auto constraint_y_n_blocks = static_cast<unsigned int>(
           std::ceil(static_cast<double>(constraint_n_blocks) /
-                    static_cast<double>(constraint_x_n_blocks));
+                    static_cast<double>(constraint_x_n_blocks)));
 
         constraint_grid_dim =
           dim3(constraint_x_n_blocks, constraint_y_n_blocks);

diff --git a/source/grid/tria.cc b/source/grid/tria.cc
index 90d1c18c762eb29ff21744913d102a04991bb908..5490a7160950977f9e9d783b13d00a6b7d84990c 100644 (file)
--- a/source/grid/tria.cc
+++ b/source/grid/tria.cc
@@ -1108,6 +1108,7 @@ namespace internal
       Assert(n_quads + 2 * n_unused_pairs + n_unused_singles ==
                tria_faces.quads.used.size(),
              ExcInternalError());
+      (void)n_quads;
 
       // how many single quads are needed in addition to n_unused_quads?
       const int additional_single_quads = new_quads_single - n_unused_singles;
@@ -1330,6 +1331,7 @@ namespace internal
           Assert(n_objects + 2 * n_unused_pairs + n_unused_singles ==
                    tria_objects.used.size(),
                  ExcInternalError());
+          (void)n_objects;
 
           // how many single objects are needed in addition to
           // n_unused_objects?
@@ -2618,6 +2620,7 @@ namespace internal
         // make sure that all subcelldata entries have been processed
         // TODO: this is not guaranteed, why?
         // AssertDimension(counter, boundary_objects_in.size());
+        (void)counter;
       }
 
 

diff --git a/source/lac/cuda_precondition.cc b/source/lac/cuda_precondition.cc
index 8573fbb16eb09401c9a0c4ba05b13c27ff047c91..8f400cfbd33c8180233db1ea20109e44fd4d64db 100644 (file)
--- a/source/lac/cuda_precondition.cc
+++ b/source/lac/cuda_precondition.cc
@@ -28,16 +28,16 @@ namespace
    */
   template <typename Number>
   cusparseStatus_t
-  cusparseXcsrilu02(cusparseHandle_t         handle,
-                    int                      m,
-                    int                      nnz,
-                    const cusparseMatDescr_t descrA,
-                    Number *                 csrValA_valM,
-                    const int *              csrRowPtrA,
-                    const int *              csrColIndA,
-                    csrilu02Info_t           info,
-                    cusparseSolvePolicy_t    policy,
-                    void *                   pBuffer)
+  cusparseXcsrilu02(cusparseHandle_t /*handle*/,
+                    int /*m*/,
+                    int /*nnz*/,
+                    const cusparseMatDescr_t /*descrA*/,
+                    Number * /*csrValA_valM*/,
+                    const int * /*csrRowPtrA*/,
+                    const int * /*csrColIndA*/,
+                    csrilu02Info_t /*info*/,
+                    cusparseSolvePolicy_t /*policy*/,
+                    void * /*pBuffer*/)
   {
     AssertThrow(false, ExcNotImplemented());
     return CUSPARSE_STATUS_INVALID_VALUE;
@@ -1289,9 +1289,13 @@ namespace CUDAWrappers
     n_nonzero_elements = A.n_nonzero_elements();
     AssertDimension(A.m(), A.n());
 
-    matrix_pointer                      = &A;
-    const auto          cusparse_matrix = A.get_cusparse_matrix();
-    const Number *const A_val_dev       = std::get<0>(cusparse_matrix);
+    matrix_pointer = &A;
+    const Number *A_val_dev;
+    std::tie(A_val_dev,
+             P_column_index_dev,
+             P_row_ptr_dev,
+             std::ignore,
+             std::ignore) = A.get_cusparse_matrix();
 
     // create a copy of the matrix entries since the algorithm works in-place.
     P_val_dev.reset(
@@ -1300,10 +1304,7 @@ namespace CUDAWrappers
                                          A_val_dev,
                                          n_nonzero_elements * sizeof(Number),
                                          cudaMemcpyDeviceToDevice);
-
-    P_column_index_dev                 = std::get<1>(cusparse_matrix);
-    P_row_ptr_dev                      = std::get<2>(cusparse_matrix);
-    const cusparseMatDescr_t mat_descr = std::get<3>(cusparse_matrix);
+    AssertCuda(cuda_status);
 
     // initialize an internal buffer we need later on
     tmp_dev.reset(Utilities::CUDA::allocate_device_data<Number>(n_rows));
@@ -1611,8 +1612,12 @@ namespace CUDAWrappers
     n_nonzero_elements = A.n_nonzero_elements();
     AssertDimension(A.m(), A.n());
 
-    const auto          cusparse_matrix = A.get_cusparse_matrix();
-    const Number *const A_val_dev       = std::get<0>(cusparse_matrix);
+    const Number *A_val_dev;
+    std::tie(A_val_dev,
+             P_column_index_dev,
+             P_row_ptr_dev,
+             std::ignore,
+             std::ignore) = A.get_cusparse_matrix();
 
     // create a copy of the matrix entries since the algorithm works in-place.
     P_val_dev.reset(
@@ -1621,10 +1626,7 @@ namespace CUDAWrappers
                                          A_val_dev,
                                          n_nonzero_elements * sizeof(Number),
                                          cudaMemcpyDeviceToDevice);
-
-    P_column_index_dev                 = std::get<1>(cusparse_matrix);
-    P_row_ptr_dev                      = std::get<2>(cusparse_matrix);
-    const cusparseMatDescr_t mat_descr = std::get<3>(cusparse_matrix);
+    AssertCuda(cuda_status);
 
     // initialize an internal buffer we need later on
     tmp_dev.reset(Utilities::CUDA::allocate_device_data<Number>(n_rows));

diff --git a/source/lac/cuda_sparse_matrix.cc b/source/lac/cuda_sparse_matrix.cc
index 5909297bf11136f68d657bde9d14c2a9c437f218..d1cca009f8bed4e07fcc21729e2b7cf92d86b6c6 100644 (file)
--- a/source/lac/cuda_sparse_matrix.cc
+++ b/source/lac/cuda_sparse_matrix.cc
@@ -265,9 +265,9 @@ namespace CUDAWrappers
     __global__ void
     linfty_norm(const typename SparseMatrix<Number>::size_type n_rows,
                 const Number *                                 val_dev,
-                const int *                                    column_index_dev,
-                const int *                                    row_ptr_dev,
-                Number *                                       sums)
+                const int * /*column_index_dev*/,
+                const int *row_ptr_dev,
+                Number *   sums)
     {
       const typename SparseMatrix<Number>::size_type row =
         threadIdx.x + blockIdx.x * blockDim.x;
@@ -668,6 +668,7 @@ namespace CUDAWrappers
                                         val_dev.get(),
                                         nnz * sizeof(Number),
                                         cudaMemcpyDeviceToDevice);
+    AssertCuda(cuda_error);
 
     return matrix_values.l2_norm();
   }

diff --git a/source/numerics/solution_transfer.cc b/source/numerics/solution_transfer.cc
index 6803c0f98a9ccdc8581a773fa8c1b7cda892872a..a22805f9eda16f92c7e22cf415624eed667f0720 100644 (file)
--- a/source/numerics/solution_transfer.cc
+++ b/source/numerics/solution_transfer.cc
@@ -332,6 +332,7 @@ SolutionTransfer<dim, VectorType, spacedim>::
     if (!cell->is_active() && cell->child(0)->coarsen_flag_set())
       ++n_coarsen_fathers;
   Assert(n_cells_to_coarsen >= 2 * n_coarsen_fathers, ExcInternalError());
+  (void)n_cells_to_coarsen;
 
   // allocate the needed memory. initialize
   // the following arrays in an efficient