Change API of a few functions cuda_kernel for consistency

diff --git a/include/deal.II/base/partitioner.templates.h b/include/deal.II/base/partitioner.templates.h
index 7ccedad7379a83434efd46f705dd34cec4c5e757..582034a7bc354319f2b474c86ccb471dffb0066d 100644 (file)
--- a/include/deal.II/base/partitioner.templates.h
+++ b/include/deal.II/base/partitioner.templates.h
@@ -161,8 +161,8 @@ namespace Utilities
               ::dealii::LinearAlgebra::CUDAWrappers::kernel::
                 gather<<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
                   temp_array_ptr,
-                  locally_owned_array.data(),
                   import_indices_plain_dev[i].first.get(),
+                  locally_owned_array.data(),
                   import_indices_plain_dev[i].second);
             }
           else

diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h
index 20c5c361ec33a98c2fed9783dc101541d7fcd891..963a9461702a8b929fb4b9cb1ef35f207df738ef 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.h
+++ b/include/deal.II/lac/cuda_kernels.h
@@ -451,12 +451,12 @@ namespace LinearAlgebra
        *
        * @ingroup CUDAWrappers
        */
-      template <typename Number>
+      template <typename Number, typename IndexType>
       __global__ void
-      set_permutated(Number *         val,
+      set_permutated(const IndexType *indices,
+                     Number *         val,
                      const Number *   v,
-                     const size_type *indices,
-                     const size_type  N);
+                     const IndexType  N);
 
 
 
@@ -469,8 +469,8 @@ namespace LinearAlgebra
       template <typename Number, typename IndexType>
       __global__ void
       gather(Number *         val,
-             const Number *   v,
              const IndexType *indices,
+             const Number *   v,
              const IndexType  N);
 
 
@@ -483,9 +483,9 @@ namespace LinearAlgebra
        */
       template <typename Number>
       __global__ void
-      add_permutated(Number *         val,
+      add_permutated(const size_type *indices,
+                     Number *         val,
                      const Number *   v,
-                     const size_type *indices,
                      const size_type  N);
     } // namespace kernel
   }   // namespace CUDAWrappers

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index 5d196ddc5cfc6e26e410f8fb83eb7660a47142f3..cd522131ba9c1c4bb822b0511e3196a09dc6c11f 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -533,12 +533,12 @@ namespace LinearAlgebra
 
 
 
-      template <typename Number>
+      template <typename Number, typename IndexType>
       __global__ void
-      set_permutated(Number *         val,
+      set_permutated(const IndexType *indices,
+                     Number *         val,
                      const Number *   v,
-                     const size_type *indices,
-                     const size_type  N)
+                     const IndexType  N)
       {
         const size_type idx_base =
           threadIdx.x + blockIdx.x * (blockDim.x * chunk_size);
@@ -555,8 +555,8 @@ namespace LinearAlgebra
       template <typename Number, typename IndexType>
       __global__ void
       gather(Number *         val,
-             const Number *   v,
              const IndexType *indices,
+             const Number *   v,
              const IndexType  N)
       {
         const IndexType idx_base =
@@ -573,9 +573,9 @@ namespace LinearAlgebra
 
       template <typename Number>
       __global__ void
-      add_permutated(Number *         val,
+      add_permutated(const size_type *indices,
+                     Number *         val,
                      const Number *   v,
-                     const size_type *indices,
                      const size_type  N)
       {
         const size_type idx_base =

diff --git a/include/deal.II/lac/la_parallel_vector.templates.h b/include/deal.II/lac/la_parallel_vector.templates.h
index ae0ec7c9612a3a217ad3f3a4f0601ab45f0c2818..b2e890c30ee504deca0ad308f57166b7fa76e9b6 100644 (file)
--- a/include/deal.II/lac/la_parallel_vector.templates.h
+++ b/include/deal.II/lac/la_parallel_vector.templates.h
@@ -301,7 +301,7 @@ namespace LinearAlgebra
                                     ::dealii::CUDAWrappers::block_size);
           ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<Number>
             <<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
-              tmp_vector.begin(), V_dev, indices_dev, n_elements);
+              indices_dev, tmp_vector.begin(), V_dev, n_elements);
 
           tmp_vector.compress(operation);
 
@@ -319,16 +319,16 @@ namespace LinearAlgebra
           if (operation == VectorOperation::add)
             ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated<
               Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
+              indices_dev,
               data.values_dev.get(),
               tmp_vector.begin(),
-              indices_dev,
               tmp_n_elements);
           else
             ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated<
               Number><<<n_blocks, ::dealii::CUDAWrappers::block_size>>>(
+              indices_dev,
               data.values_dev.get(),
               tmp_vector.begin(),
-              indices_dev,
               tmp_n_elements);
 
           ::dealii::Utilities::CUDA::free(indices_dev);

diff --git a/source/lac/cuda_kernels.cu b/source/lac/cuda_kernels.cu
index 317399c74d427632c5ec29d140d159813f1e4c0e..a1c83d59352d4de6c4cd8d799072b84b8932334f 100644 (file)
--- a/source/lac/cuda_kernels.cu
+++ b/source/lac/cuda_kernels.cu
@@ -94,12 +94,12 @@ namespace LinearAlgebra
                  const float *   V_val,
                  const size_type N);
       template __global__ void
-      equ(float *         val,
-          const float     a,
-          const float *   V_val,
-          const float     b,
-          const float *   W_val,
-          const size_type N);
+      equ<float>(float *         val,
+                 const float     a,
+                 const float *   V_val,
+                 const float     b,
+                 const float *   W_val,
+                 const size_type N);
       template __global__ void
       add_and_dot<float>(float *         res,
                          float *         v1,
@@ -110,19 +110,19 @@ namespace LinearAlgebra
       template __global__ void
       set<float>(float *val, const float s, const size_type N);
       template __global__ void
-      set_permutated<float>(float *          val,
-                            const float *    v,
-                            const size_type *indices,
-                            const size_type  N);
+      set_permutated<float, size_type>(const size_type *indices,
+                                       float *          val,
+                                       const float *    v,
+                                       const size_type  N);
       template __global__ void
-      gather(float *          val,
-             const float *    v,
-             const size_type *indices,
-             const size_type  N);
+      gather<float, size_type>(float *          val,
+                               const size_type *indices,
+                               const float *    v,
+                               const size_type  N);
       template __global__ void
-      add_permutated<float>(float *          val,
+      add_permutated<float>(const size_type *indices,
+                            float *          val,
                             const float *    v,
-                            const size_type *indices,
                             const size_type  N);
 
 
@@ -194,12 +194,12 @@ namespace LinearAlgebra
                   const double *  V_val,
                   const size_type N);
       template __global__ void
-      equ(double *        val,
-          const double    a,
-          const double *  V_val,
-          const double    b,
-          const double *  W_val,
-          const size_type N);
+      equ<double>(double *        val,
+                  const double    a,
+                  const double *  V_val,
+                  const double    b,
+                  const double *  W_val,
+                  const size_type N);
       template __global__ void
       add_and_dot<double>(double *        res,
                           double *        v1,
@@ -210,19 +210,19 @@ namespace LinearAlgebra
       template __global__ void
       set<double>(double *val, const double s, const size_type N);
       template __global__ void
-      set_permutated<double>(double *         val,
-                             const double *   v,
-                             const size_type *indices,
-                             const size_type  N);
+      set_permutated<double, size_type>(const size_type *indices,
+                                        double *         val,
+                                        const double *   v,
+                                        const size_type  N);
       template __global__ void
-      gather(double *         val,
-             const double *   v,
-             const size_type *indices,
-             const size_type  N);
+      gather<double, size_type>(double *         val,
+                                const size_type *indices,
+                                const double *   v,
+                                const size_type  N);
       template __global__ void
-      add_permutated<double>(double *         val,
+      add_permutated<double>(const size_type *indices,
+                             double *         val,
                              const double *   v,
-                             const size_type *indices,
                              const size_type  N);
     } // namespace kernel
   }   // namespace CUDAWrappers