From: Bruno Turcksin Date: Fri, 8 Feb 2019 14:36:02 +0000 (+0000) Subject: Change API of a few functions cuda_kernel for consistency X-Git-Tag: v9.1.0-rc1~349^2~3 X-Git-Url: https://gitweb.dealii.org/cgi-bin/gitweb.cgi?a=commitdiff_plain;h=42021d3a02ff37efd241e0b77b1ebd05bb01daa8;p=dealii.git Change API of a few functions cuda_kernel for consistency --- diff --git a/include/deal.II/base/partitioner.templates.h b/include/deal.II/base/partitioner.templates.h index 7ccedad737..582034a7bc 100644 --- a/include/deal.II/base/partitioner.templates.h +++ b/include/deal.II/base/partitioner.templates.h @@ -161,8 +161,8 @@ namespace Utilities ::dealii::LinearAlgebra::CUDAWrappers::kernel:: gather<<>>( temp_array_ptr, - locally_owned_array.data(), import_indices_plain_dev[i].first.get(), + locally_owned_array.data(), import_indices_plain_dev[i].second); } else diff --git a/include/deal.II/lac/cuda_kernels.h b/include/deal.II/lac/cuda_kernels.h index 20c5c361ec..963a946170 100644 --- a/include/deal.II/lac/cuda_kernels.h +++ b/include/deal.II/lac/cuda_kernels.h @@ -451,12 +451,12 @@ namespace LinearAlgebra * * @ingroup CUDAWrappers */ - template + template __global__ void - set_permutated(Number * val, + set_permutated(const IndexType *indices, + Number * val, const Number * v, - const size_type *indices, - const size_type N); + const IndexType N); @@ -469,8 +469,8 @@ namespace LinearAlgebra template __global__ void gather(Number * val, - const Number * v, const IndexType *indices, + const Number * v, const IndexType N); @@ -483,9 +483,9 @@ namespace LinearAlgebra */ template __global__ void - add_permutated(Number * val, + add_permutated(const size_type *indices, + Number * val, const Number * v, - const size_type *indices, const size_type N); } // namespace kernel } // namespace CUDAWrappers diff --git a/include/deal.II/lac/cuda_kernels.templates.h b/include/deal.II/lac/cuda_kernels.templates.h index 5d196ddc5c..cd522131ba 100644 --- a/include/deal.II/lac/cuda_kernels.templates.h +++ b/include/deal.II/lac/cuda_kernels.templates.h @@ -533,12 +533,12 @@ namespace LinearAlgebra - template + template __global__ void - set_permutated(Number * val, + set_permutated(const IndexType *indices, + Number * val, const Number * v, - const size_type *indices, - const size_type N) + const IndexType N) { const size_type idx_base = threadIdx.x + blockIdx.x * (blockDim.x * chunk_size); @@ -555,8 +555,8 @@ namespace LinearAlgebra template __global__ void gather(Number * val, - const Number * v, const IndexType *indices, + const Number * v, const IndexType N) { const IndexType idx_base = @@ -573,9 +573,9 @@ namespace LinearAlgebra template __global__ void - add_permutated(Number * val, + add_permutated(const size_type *indices, + Number * val, const Number * v, - const size_type *indices, const size_type N) { const size_type idx_base = diff --git a/include/deal.II/lac/la_parallel_vector.templates.h b/include/deal.II/lac/la_parallel_vector.templates.h index ae0ec7c961..b2e890c30e 100644 --- a/include/deal.II/lac/la_parallel_vector.templates.h +++ b/include/deal.II/lac/la_parallel_vector.templates.h @@ -301,7 +301,7 @@ namespace LinearAlgebra ::dealii::CUDAWrappers::block_size); ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated <<>>( - tmp_vector.begin(), V_dev, indices_dev, n_elements); + indices_dev, tmp_vector.begin(), V_dev, n_elements); tmp_vector.compress(operation); @@ -319,16 +319,16 @@ namespace LinearAlgebra if (operation == VectorOperation::add) ::dealii::LinearAlgebra::CUDAWrappers::kernel::add_permutated< Number><<>>( + indices_dev, data.values_dev.get(), tmp_vector.begin(), - indices_dev, tmp_n_elements); else ::dealii::LinearAlgebra::CUDAWrappers::kernel::set_permutated< Number><<>>( + indices_dev, data.values_dev.get(), tmp_vector.begin(), - indices_dev, tmp_n_elements); ::dealii::Utilities::CUDA::free(indices_dev); diff --git a/source/lac/cuda_kernels.cu b/source/lac/cuda_kernels.cu index 317399c74d..a1c83d5935 100644 --- a/source/lac/cuda_kernels.cu +++ b/source/lac/cuda_kernels.cu @@ -94,12 +94,12 @@ namespace LinearAlgebra const float * V_val, const size_type N); template __global__ void - equ(float * val, - const float a, - const float * V_val, - const float b, - const float * W_val, - const size_type N); + equ(float * val, + const float a, + const float * V_val, + const float b, + const float * W_val, + const size_type N); template __global__ void add_and_dot(float * res, float * v1, @@ -110,19 +110,19 @@ namespace LinearAlgebra template __global__ void set(float *val, const float s, const size_type N); template __global__ void - set_permutated(float * val, - const float * v, - const size_type *indices, - const size_type N); + set_permutated(const size_type *indices, + float * val, + const float * v, + const size_type N); template __global__ void - gather(float * val, - const float * v, - const size_type *indices, - const size_type N); + gather(float * val, + const size_type *indices, + const float * v, + const size_type N); template __global__ void - add_permutated(float * val, + add_permutated(const size_type *indices, + float * val, const float * v, - const size_type *indices, const size_type N); @@ -194,12 +194,12 @@ namespace LinearAlgebra const double * V_val, const size_type N); template __global__ void - equ(double * val, - const double a, - const double * V_val, - const double b, - const double * W_val, - const size_type N); + equ(double * val, + const double a, + const double * V_val, + const double b, + const double * W_val, + const size_type N); template __global__ void add_and_dot(double * res, double * v1, @@ -210,19 +210,19 @@ namespace LinearAlgebra template __global__ void set(double *val, const double s, const size_type N); template __global__ void - set_permutated(double * val, - const double * v, - const size_type *indices, - const size_type N); + set_permutated(const size_type *indices, + double * val, + const double * v, + const size_type N); template __global__ void - gather(double * val, - const double * v, - const size_type *indices, - const size_type N); + gather(double * val, + const size_type *indices, + const double * v, + const size_type N); template __global__ void - add_permutated(double * val, + add_permutated(const size_type *indices, + double * val, const double * v, - const size_type *indices, const size_type N); } // namespace kernel } // namespace CUDAWrappers