From: Bruno Turcksin <bruno.turcksin@gmail.com>
Date: Tue, 18 Dec 2018 14:14:09 +0000 (+0000)
Subject: Make gather kernel more general
X-Git-Tag: v9.1.0-rc1~464^2~1
X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=87b7a2fdc071011142e088b1db8fbceb82fd0698;p=dealii.git

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 89e9420afc..e953269398 100644
--- 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 e6ca63f88d..5d196ddc5c 100644
--- 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 49230c517c..317399c74d 100644
--- 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,