]> https://gitweb.dealii.org/ - dealii.git/commitdiff
Add gather cuda kernel
authorBruno Turcksin <bruno.turcksin@gmail.com>
Mon, 26 Nov 2018 13:48:58 +0000 (13:48 +0000)
committerDaniel Arndt <daniel.arndt@iwr.uni-heidelberg.de>
Fri, 4 Jan 2019 14:14:38 +0000 (15:14 +0100)
include/deal.II/lac/cuda_kernels.h
include/deal.II/lac/cuda_kernels.templates.h

index 6ce0f431f3305fbfa0d481c1ecf8e7c1fd677a5f..89e9420afc70af251bcbf0b064fe722554da20f4 100644 (file)
@@ -460,6 +460,21 @@ namespace LinearAlgebra
 
 
 
+      /**
+       * Set each element @v val to @p v using @p indices as permutation, i.e.,
+       * <tt>val[i] = v[indices[i]]</tt>.
+       *
+       * @ingroup CUDAWrappers
+       */
+      template <typename Number>
+      __global__ void
+      gather(Number *         val,
+             const Number *   v,
+             const size_type *indices,
+             const size_type  N);
+
+
+
       /**
        * Add each element @v val to @p v using @p indices as permutation, i.e.,
        * <tt>val[indices[i]] += v[i]</tt>.
index 7a8185f302e22eba946dfda5e9e738102fc786d5..e6ca63f88dfd18e32d9008faab5b9627b2093f00 100644 (file)
@@ -552,6 +552,25 @@ namespace LinearAlgebra
 
 
 
+      template <typename Number>
+      __global__ void
+      gather(Number *         val,
+             const Number *   v,
+             const size_type *indices,
+             const size_type  N)
+      {
+        const size_type 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;
+            if (idx < N)
+              val[idx] = v[indices[idx]];
+          }
+      }
+
+
+
       template <typename Number>
       __global__ void
       add_permutated(Number *         val,

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.