From b7a3464775fb2cdb7e40f18c6cd8ab8c725db192 Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Mon, 26 Nov 2018 13:48:58 +0000 Subject: [PATCH] Add gather cuda kernel --- include/deal.II/lac/cuda_kernels.h | 15 +++++++++++++++ include/deal.II/lac/cuda_kernels.templates.h | 19 +++++++++++++++++++ 2 files changed, 34 insertions(+) diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h index 6ce0f431f3..89e9420afc 100644 --- a/include/deal.II/lac/cuda_kernels.h +++ b/include/deal.II/lac/cuda_kernels.h @@ -460,6 +460,21 @@ namespace LinearAlgebra + /** + * Set each element @v val to @p v using @p indices as permutation, i.e., + * val[i] = v[indices[i]]. + * + * @ingroup CUDAWrappers + */ + template + __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., * val[indices[i]] += v[i]. diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h index 7a8185f302..e6ca63f88d 100644 --- a/include/deal.II/lac/cuda_kernels.templates.h +++ b/include/deal.II/lac/cuda_kernels.templates.h @@ -552,6 +552,25 @@ namespace LinearAlgebra + template + __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 __global__ void add_permutated(Number * val, -- 2.39.5