From 1624c99f40004ff3e5772a416d2863dbb04a4225 Mon Sep 17 00:00:00 2001 From: Bruno Turcksin Date: Mon, 31 Oct 2016 15:15:09 -0400 Subject: [PATCH] Add mean_value to cuda_vector. --- include/deal.II/lac/cuda_vector.h | 6 +++ include/deal.II/lac/la_vector.h | 2 +- include/deal.II/lac/vector_space_vector.h | 2 +- source/lac/cuda_vector.cu | 56 +++++++++++++++++++++++ tests/cuda/cuda_vector_02.cc | 2 + 5 files changed, 66 insertions(+), 2 deletions(-) diff --git a/include/deal.II/lac/cuda_vector.h b/include/deal.II/lac/cuda_vector.h index 8a925f6346..c91751d486 100644 --- a/include/deal.II/lac/cuda_vector.h +++ b/include/deal.II/lac/cuda_vector.h @@ -44,6 +44,7 @@ namespace LinearAlgebra class Vector: public VectorSpaceVector { public: + typedef typename VectorSpaceVector::value_type value_type; typedef typename VectorSpaceVector::size_type size_type; typedef typename VectorSpaceVector::real_type real_type; @@ -155,6 +156,11 @@ namespace LinearAlgebra */ virtual void equ(const Number a, const VectorSpaceVector &V) override; + /** + * Return the mean value of all the entries of this vector. + */ + virtual value_type mean_value() const override; + /** * Return the l1 norm of the vector (i.e., the sum of the * absolute values of all entries among all processors). diff --git a/include/deal.II/lac/la_vector.h b/include/deal.II/lac/la_vector.h index 4d6b983039..144315e5c0 100644 --- a/include/deal.II/lac/la_vector.h +++ b/include/deal.II/lac/la_vector.h @@ -179,7 +179,7 @@ namespace LinearAlgebra virtual void equ(const Number a, const VectorSpaceVector &V); /** - * Return the mean value of the element of this vector. + * Return the mean value of all the entries of this vector. */ virtual typename VectorSpaceVector::value_type mean_value() const; diff --git a/include/deal.II/lac/vector_space_vector.h b/include/deal.II/lac/vector_space_vector.h index 491104467f..cb8c95fdc8 100644 --- a/include/deal.II/lac/vector_space_vector.h +++ b/include/deal.II/lac/vector_space_vector.h @@ -129,7 +129,7 @@ namespace LinearAlgebra virtual void equ(const Number a, const VectorSpaceVector &V) = 0; /** - * Return the mean value of the element of this vector. + * Return the mean value of all the entries of this vector. */ virtual value_type mean_value() const = 0; diff --git a/source/lac/cuda_vector.cu b/source/lac/cuda_vector.cu index 1e3b295edf..c0aa183792 100644 --- a/source/lac/cuda_vector.cu +++ b/source/lac/cuda_vector.cu @@ -93,6 +93,33 @@ namespace LinearAlgebra } + + template + struct ElemSum + { + __device__ static Number reduction_op(const Number a, const Number b) + { + return (a + b); + } + + __device__ static Number atomic_op(Number *dst, const Number a) + { + return atomicAdd_wrapper(dst, a); + } + + __device__ static Number element_wise_op(const Number a) + { + return a; + } + + __device__ static Number null_value() + { + return Number(); + } + }; + + + template struct L1Norm { @@ -118,6 +145,7 @@ namespace LinearAlgebra }; + template struct LInfty { @@ -864,6 +892,34 @@ namespace LinearAlgebra + template + typename Vector::value_type Vector::mean_value() const + { + Number *result_device; + cudaError_t error_code = cudaMalloc(&result_device, sizeof(Number)); + AssertCuda(error_code); + error_code = cudaMemset(result_device, Number(), sizeof(Number)); + + const int n_blocks = 1 + (n_elements-1)/(CHUNK_SIZE*BLOCK_SIZE); + internal::reduction> + <<>> ( + result_device, val, + n_elements); + + // Copy the result back to the host + Number result; + error_code = cudaMemcpy(&result, result_device, sizeof(Number), + cudaMemcpyDeviceToHost); + AssertCuda(error_code); + // Free the memory on the device + error_code = cudaFree(result_device); + AssertCuda(error_code); + + return result/static_cast::value_type>(n_elements); + } + + + template typename Vector::real_type Vector::l1_norm() const { diff --git a/tests/cuda/cuda_vector_02.cc b/tests/cuda/cuda_vector_02.cc index c17999c603..09b317cfc0 100644 --- a/tests/cuda/cuda_vector_02.cc +++ b/tests/cuda/cuda_vector_02.cc @@ -84,6 +84,8 @@ void test() AssertThrow(2.*read_write_2[i]==read_write_3[i], ExcMessage("Problem in equ.")); + AssertThrow(b.mean_value()==54.50, ExcMessage("Problem in mean_value.")); + AssertThrow(b.l1_norm()==5450., ExcMessage("Problem in l1_norm.")); const double eps=1e-3; -- 2.39.5