]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Make gather kernel more general
authorBruno Turcksin <bruno.turcksin@gmail.com>
Tue, 18 Dec 2018 14:14:09 +0000 (14:14 +0000)
committerDaniel Arndt <daniel.arndt@iwr.uni-heidelberg.de>
Fri, 4 Jan 2019 14:14:39 +0000 (15:14 +0100)
This make the index type a template parameter so that gather
can be used export_to_ghosted_array_start when using 64 bit indices.

include/deal.II/lac/cuda_kernels.h
include/deal.II/lac/cuda_kernels.templates.h
source/lac/cuda_kernels.cu

index 89e9420afc70af251bcbf0b064fe722554da20f4..e953269398520aa8c5f342b621ee26af526aabd0 100644 (file)
@@ -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);
 
 
 
index e6ca63f88dfd18e32d9008faab5b9627b2093f00..5d196ddc5cfc6e26e410f8fb83eb7660a47142f3 100644 (file)
@@ -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]];
           }
index 49230c517cc3339e239c302e79c2be86233c1b72..317399c74d427632c5ec29d140d159813f1e4c0e 100644 (file)
@@ -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,

In the beginning the Universe was created. This has made a lot of people very angry and has been widely regarded as a bad move.

Douglas Adams


Typeset in Trocchi and Trocchi Bold Sans Serif.