Make gather kernel more general

This make the index type a template parameter so that gather
can be used export_to_ghosted_array_start when using 64 bit indices.

diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h
index 89e9420afc70af251bcbf0b064fe722554da20f4..e953269398520aa8c5f342b621ee26af526aabd0 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.h
+++ b/include/deal.II/lac/cuda_kernels.h
@@ -466,12 +466,12 @@ namespace LinearAlgebra
        *
        * @ingroup CUDAWrappers
        */
-      template <typename Number>
+      template <typename Number, typename IndexType>
       __global__ void
       gather(Number *         val,
              const Number *   v,
-             const size_type *indices,
-             const size_type  N);
+             const IndexType *indices,
+             const IndexType  N);
 
 
 

diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h
index e6ca63f88dfd18e32d9008faab5b9627b2093f00..5d196ddc5cfc6e26e410f8fb83eb7660a47142f3 100644 (file)
--- a/include/deal.II/lac/cuda_kernels.templates.h
+++ b/include/deal.II/lac/cuda_kernels.templates.h
@@ -552,18 +552,18 @@ namespace LinearAlgebra
 
 
 
-      template <typename Number>
+      template <typename Number, typename IndexType>
       __global__ void
       gather(Number *         val,
              const Number *   v,
-             const size_type *indices,
-             const size_type  N)
+             const IndexType *indices,
+             const IndexType  N)
       {
-        const size_type idx_base =
+        const IndexType idx_base =
           threadIdx.x + blockIdx.x * (blockDim.x * chunk_size);
         for (unsigned int i = 0; i < chunk_size; ++i)
           {
-            const size_type idx = idx_base + i * block_size;
+            const IndexType idx = idx_base + i * block_size;
             if (idx < N)
               val[idx] = v[indices[idx]];
           }

diff --git a/source/lac/cuda_kernels.cu b/source/lac/cuda_kernels.cu
index 49230c517cc3339e239c302e79c2be86233c1b72..317399c74d427632c5ec29d140d159813f1e4c0e 100644 (file)
--- a/source/lac/cuda_kernels.cu
+++ b/source/lac/cuda_kernels.cu
@@ -115,6 +115,11 @@ namespace LinearAlgebra
                             const size_type *indices,
                             const size_type  N);
       template __global__ void
+      gather(float *          val,
+             const float *    v,
+             const size_type *indices,
+             const size_type  N);
+      template __global__ void
       add_permutated<float>(float *          val,
                             const float *    v,
                             const size_type *indices,
@@ -210,6 +215,11 @@ namespace LinearAlgebra
                              const size_type *indices,
                              const size_type  N);
       template __global__ void
+      gather(double *         val,
+             const double *   v,
+             const size_type *indices,
+             const size_type  N);
+      template __global__ void
       add_permutated<double>(double *         val,
                              const double *   v,
                              const size_type *indices,