From 87b7a2fdc071011142e088b1db8fbceb82fd0698 Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Tue, 18 Dec 2018 14:14:09 +0000 Subject: [PATCH] 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. --- include/deal.II/lac/cuda_kernels.h | 6 +++--- include/deal.II/lac/cuda_kernels.templates.h | 10 +++++----- source/lac/cuda_kernels.cu | 10 ++++++++++ 3 files changed, 18 insertions(+), 8 deletions(-) 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 + template __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 + template __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 * 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 * val, const double * v, const size_type *indices, -- 2.39.5