diff --git a/src/TNL/Containers/Expressions/ExpressionTemplates.h b/src/TNL/Containers/Expressions/ExpressionTemplates.h index 8b82f5422076385f117e9b9f00136614861552dd..f942ceb8cc0fa05772725b90d0b9b4323c6b678b 100644 --- a/src/TNL/Containers/Expressions/ExpressionTemplates.h +++ b/src/TNL/Containers/Expressions/ExpressionTemplates.h @@ -2116,6 +2116,89 @@ operator,( const Containers::Expressions::BinaryExpressionTemplate< L1, L2, LOpe return TNL::sum( a * b ); } +template< typename T1, + typename T2, + template< typename, typename > class Operation > +auto +dot( const Containers::Expressions::BinaryExpressionTemplate< T1, T2, Operation >& a, + const typename Containers::Expressions::BinaryExpressionTemplate< T1, T2, Operation >::RealType& b ) +-> decltype( TNL::sum( a * b ) ) +{ + return TNL::sum( a * b ); +} + +template< typename L1, + template< typename > class LOperation, + typename R1, + typename R2, + template< typename, typename > class ROperation > +auto +dot( const Containers::Expressions::UnaryExpressionTemplate< L1, LOperation >& a, + const typename Containers::Expressions::BinaryExpressionTemplate< R1, R2, ROperation >& b ) +-> decltype( TNL::sum( a * b ) ) +{ + return TNL::sum( a * b ); +} + +template< typename L1, + typename L2, + template< typename, typename > class LOperation, + typename R1, + template< typename > class ROperation > +auto +dot( const Containers::Expressions::BinaryExpressionTemplate< L1, L2, LOperation >& a, + const typename Containers::Expressions::UnaryExpressionTemplate< R1,ROperation >& b ) +-> decltype( TNL::sum( a * b ) ) +{ + return TNL::sum( a * b ); +} + + +//// +// Evaluation with reduction +template< typename Vector, + typename T1, + typename T2, + template< typename, typename > class Operation, + typename Reduction, + typename VolatileReduction, + typename Result > +Result evaluateAndReduce( Vector& lhs, + const Containers::Expressions::BinaryExpressionTemplate< T1, T2, Operation >& expression, + Reduction& reduction, + VolatileReduction& volatileReduction, + const Result& zero ) +{ + using RealType = typename Vector::RealType; + using IndexType = typename Vector::IndexType; + using DeviceType = typename Vector::DeviceType; + + RealType* lhs_data = lhs.getData(); + auto fetch = [=] __cuda_callable__ ( IndexType i ) -> RealType { return ( lhs_data[ i ] = expression[ i ] ); }; + return Containers::Algorithms::Reduction< DeviceType >::reduce( lhs.getSize(), reduction, volatileReduction, fetch, zero ); +} + +template< typename Vector, + typename T1, + template< typename > class Operation, + typename Reduction, + typename VolatileReduction, + typename Result > +Result evaluateAndReduce( Vector& lhs, + const Containers::Expressions::UnaryExpressionTemplate< T1, Operation >& expression, + Reduction& reduction, + VolatileReduction& volatileReduction, + const Result& zero ) +{ + using RealType = typename Vector::RealType; + using IndexType = typename Vector::IndexType; + using DeviceType = typename Vector::DeviceType; + + RealType* lhs_data = lhs.getData(); + auto fetch = [=] __cuda_callable__ ( IndexType i ) -> RealType { return ( lhs_data[ i ] = expression[ i ] ); }; + return Containers::Algorithms::Reduction< DeviceType >::reduce( lhs.getSize(), reduction, volatileReduction, fetch, zero ); +} + //// // Output stream template< typename T1, diff --git a/src/TNL/Containers/Expressions/StaticExpressionTemplates.h b/src/TNL/Containers/Expressions/StaticExpressionTemplates.h index bd04a686d3ced76e288d96be5acb3687f431f2ab..df045bbba25e15decebbdc0288c311028d4d24b0 100644 --- a/src/TNL/Containers/Expressions/StaticExpressionTemplates.h +++ b/src/TNL/Containers/Expressions/StaticExpressionTemplates.h @@ -2388,5 +2388,45 @@ dot( const Containers::Expressions::StaticBinaryExpressionTemplate< L1, L2, LOpe return TNL::sum( a * b ); } +//// +// Evaluation with reduction +template< typename Vector, + typename T1, + typename T2, + template< typename, typename > class Operation, + typename Reduction, + typename VolatileReduction, + typename Result > +__cuda_callable__ +Result evaluateAndReduce( Vector& lhs, + const Containers::Expressions::StaticBinaryExpressionTemplate< T1, T2, Operation >& expression, + Reduction& reduction, + VolatileReduction& volatileReduction, + const Result& zero ) +{ + Result result( zero ); + for( int i = 0; i < Vector::getSize(); i++ ) + reduction( result, lhs[ i ] = expression[ i ] ); + return result; +} + +template< typename Vector, + typename T1, + template< typename > class Operation, + typename Reduction, + typename VolatileReduction, + typename Result > +__cuda_callable__ +Result evaluateAndReduce( Vector& lhs, + const Containers::Expressions::StaticUnaryExpressionTemplate< T1, Operation >& expression, + Reduction& reduction, + VolatileReduction& volatileReduction, + const Result& zero ) +{ + Result result( zero ); + for( int i = 0; i < Vector::getSize(); i++ ) + reduction( result, lhs[ i ] = expression[ i ] ); + return result; +} } // namespace TNL diff --git a/src/TNL/Solvers/ODE/Euler_impl.h b/src/TNL/Solvers/ODE/Euler.hpp similarity index 100% rename from src/TNL/Solvers/ODE/Euler_impl.h rename to src/TNL/Solvers/ODE/Euler.hpp diff --git a/src/TNL/TypeTraits.h b/src/TNL/TypeTraits.h new file mode 100644 index 0000000000000000000000000000000000000000..bb7e4a4d1efae8ecb096addb5c645fda378a6ae5 --- /dev/null +++ b/src/TNL/TypeTraits.h @@ -0,0 +1,19 @@ +/*************************************************************************** + TypeTraits.h - description + ------------------- + begin : Jun 25, 2019 + copyright : (C) 2019 by Tomas Oberhuber et al. + email : tomas.oberhuber@fjfi.cvut.cz + ***************************************************************************/ + +/* See Copyright Notice in tnl/Copyright */ + +namespace TNL { + +template< typename T > +struct ViewType +{ + using Type = T; +}; + +} //namespace TNL \ No newline at end of file diff --git a/src/UnitTests/Containers/VectorTest-7.h b/src/UnitTests/Containers/VectorTest-7.h index ecf6fabbf698a7102b5e9c47797380a86dee4d8a..459e546dd618cac7b2d36b031053df9c07445152 100644 --- a/src/UnitTests/Containers/VectorTest-7.h +++ b/src/UnitTests/Containers/VectorTest-7.h @@ -399,6 +399,44 @@ TYPED_TEST( VectorTest, sign ) EXPECT_NEAR( sign( u ).getElement( i ), v.getElement( i ), 1.0e-6 ); } +// NOTE: The following lambdas cannot be inside the test because of nvcc ( v. 10.1.105 ) +// error #3049-D: The enclosing parent function ("TestBody") for an extended __host__ __device__ lambda cannot have private or protected access within its class +template< typename VectorView > +typename VectorView::RealType +performEvaluateAndReduce( VectorView& u, VectorView& v, VectorView& w ) +{ + using RealType = typename VectorView::RealType; + + auto reduction = [] __cuda_callable__ ( RealType& a, const RealType& b ) { a += b; }; + auto volatileReduction = [] __cuda_callable__ ( volatile RealType& a, volatile RealType& b ) { a += b; }; + return evaluateAndReduce( w, u * v, reduction, volatileReduction, ( RealType ) 0.0 ); +} + +TYPED_TEST( VectorTest, evaluateAndReduce ) +{ + using VectorType = typename TestFixture::VectorType; + using ViewType = typename TestFixture::ViewType; + using RealType = typename VectorType::RealType; + using IndexType = typename VectorType::IndexType; + const int size = VECTOR_TEST_SIZE; + + VectorType _u( size ), _v( size ), _w( size ); + ViewType u( _u ), v( _v ), w( _w ); + RealType aux( 0.0 ); + for( int i = 0; i < size; i++ ) + { + const RealType x = i; + const RealType y = size / 2 - i; + u.setElement( i, x ); + v.setElement( i, y ); + aux += x * y; + } + auto r = performEvaluateAndReduce( u, v, w ); + EXPECT_TRUE( w == u * v ); + EXPECT_NEAR( aux, r, 1.0e-5 ); +} + + #endif // HAVE_GTEST #include "../main.h"