From f343ff947ead9220a9123e9a0399ae4af430d0bf Mon Sep 17 00:00:00 2001 From: Tomas Oberhuber <tomas.oberhuber@fjfi.cvut.cz> Date: Wed, 6 Jan 2016 10:51:13 +0100 Subject: [PATCH] Implementing operator functions. --- src/functions/tnlBoundaryOperatorFunction.h | 24 +++ src/functions/tnlMeshFunctionEvaluator.h | 67 ++++--- src/functions/tnlMeshFunctionEvaluator_impl.h | 94 ++++++++-- src/functions/tnlMeshFunction_impl.h | 41 +--- src/functions/tnlOperatorFunction.h | 27 +-- src/functions/tnlOperatorFunction_impl.h | 17 +- src/operators/tnlOperatorEnumerator.h | 175 ------------------ src/operators/tnlOperatorEnumerator_impl.h | 141 -------------- tests/unit-tests/tnlApproximationError_impl.h | 4 +- 9 files changed, 171 insertions(+), 419 deletions(-) create mode 100644 src/functions/tnlBoundaryOperatorFunction.h delete mode 100644 src/operators/tnlOperatorEnumerator.h delete mode 100644 src/operators/tnlOperatorEnumerator_impl.h diff --git a/src/functions/tnlBoundaryOperatorFunction.h b/src/functions/tnlBoundaryOperatorFunction.h new file mode 100644 index 0000000000..aa635d6c58 --- /dev/null +++ b/src/functions/tnlBoundaryOperatorFunction.h @@ -0,0 +1,24 @@ +/*************************************************************************** + tnlBoundaryOperatorFunction.h - description + ------------------- + begin : Jan 6, 2016 + copyright : (C) 2016 by oberhuber + email : tomas.oberhuber@fjfi.cvut.cz + ***************************************************************************/ + +/*************************************************************************** + * * + * This program is free software; you can redistribute it and/or modify * + * it under the terms of the GNU General Public License as published by * + * the Free Software Foundation; either version 2 of the License, or * + * (at your option) any later version. * + * * + ***************************************************************************/ + +#ifndef TNLBOUNDARYOPERATORFUNCTION_H +#define TNLBOUNDARYOPERATORFUNCTION_H + + + +#endif /* TNLBOUNDARYOPERATORFUNCTION_H */ + diff --git a/src/functions/tnlMeshFunctionEvaluator.h b/src/functions/tnlMeshFunctionEvaluator.h index 3a41818d07..b78484baa1 100644 --- a/src/functions/tnlMeshFunctionEvaluator.h +++ b/src/functions/tnlMeshFunctionEvaluator.h @@ -20,6 +20,7 @@ #include <mesh/tnlGrid.h> #include <functions/tnlMeshFunction.h> +#include <functions/tnlOperatorFunction.h> template< typename OutMeshFunction, typename InFunction > @@ -87,36 +88,54 @@ class tnlMeshFunctionEvaluator }; }; -/*template< int Dimensions, - typename MeshReal1, - typename MeshReal2, - typename MeshDevice1, - typename MeshDevice2, - typename MeshIndex1, - typename MeshIndex2, - int MeshEntityDimensions, - typename Real > -class tnlMeshFunctionEvaluator< tnlMeshFunction< tnlGrid< Dimensions, MeshReal1, MeshDevice1, MeshIndex1 >, MeshEntityDimensions, Real >, - tnlMeshFunction< tnlGrid< Dimensions, MeshReal2, MeshDevice2, MeshIndex2 >, MeshEntityDimensions, Real > > +template< typename OutMeshFunction, + typename Operator, + typename Function > +class tnlMeshFunctionEvaluator< OutMeshFunction, tnlOperatorFunction< Operator, Function > > { public: - typedef tnlGrid< Dimensions, MeshReal1, MeshDevice1, MeshIndex1 > Mesh1; - typedef tnlGrid< Dimensions, MeshReal2, MeshDevice2, MeshIndex2 > Mesh2; - typedef tnlMeshFunction< Mesh1, MeshEntityDimensions, Real > OutMeshFunction; - typedef tnlMeshFunction< Mesh2, MeshEntityDimensions, Real > InFunction; + typedef typename OutMeshFunction::MeshType MeshType; + typedef typename MeshType::RealType MeshRealType; + typedef typename MeshType::DeviceType MeshDeviceType; + typedef typename MeshType::IndexType MeshIndexType; + typedef typename OutMeshFunction::Real RealType; + typedef tnlOperatorFunction< Operator, Function > OperatorFunctionType; + + static void evaluateAllEntities( OutMeshFunction& meshFunction, + const OperatorFunctionType& function, + const RealType& time = 0.0, + const RealType& outFunctionMultiplicator = 0.0, + const RealType& inFunctionMultiplicator = 1.0 ) + { + evaluateInteriorEntities( meshFunction, function, time, outFunctionMultiplicator, inFunctionMultiplicator ); + }; + + static void evaluateInteriorEntities( OutMeshFunction& meshFunction, + const OperatorFunctionType& function, + const RealType& time = 0.0, + const RealType& outFunctionMultiplicator = 0.0, + const RealType& inFunctionMultiplicator = 1.0 ); - static void assign( OutMeshFunction& f1, - const InFunction& f2 ) + class TraverserUserData { - if( f1.getMesh().getDimensions() == f2.getMesh().getDimensions() ) - f1.getData() = f2.getData(); - else - { - //TODO: Interpolace - } + public: + TraverserUserData( const OperatorFunctionType* operatorFunction, + const RealType* time, + OutMeshFunction* meshFunction, + const RealType* outFunctionMultiplicator, + const RealType* inFunctionMultiplicator ) + : meshFunction( meshFunction ), operatorFunction( operatorFunction ),time( time ), + outFunctionMultiplicator( outFunctionMultiplicator ), + inFunctionMultiplicator( inFunctionMultiplicator ){} + + protected: + OutMeshFunction* meshFunction; + const OperatorFunctionType* operatorFunction; + const RealType *time, *outFunctionMultiplicator, *inFunctionMultiplicator; + }; -};*/ +}; #endif /* TNLMESHFUNCTIONEVALUATOR_H */ diff --git a/src/functions/tnlMeshFunctionEvaluator_impl.h b/src/functions/tnlMeshFunctionEvaluator_impl.h index 5069eb87cf..557f80d913 100644 --- a/src/functions/tnlMeshFunctionEvaluator_impl.h +++ b/src/functions/tnlMeshFunctionEvaluator_impl.h @@ -18,6 +18,9 @@ #ifndef TNLMESHFUNCTIONEVALUATOR_IMPL_H #define TNLMESHFUNCTIONEVALUATOR_IMPL_H +#include <functions/tnlMeshFunctionEvaluator.h> +#include <mesh/tnlTraverser.h> + template< typename OutMeshFunction, typename InFunction > void @@ -63,12 +66,12 @@ template< typename OutMeshFunction, typename InFunction > void tnlMeshFunctionEvaluator< OutMeshFunction, InFunction >:: -assign( OutMeshFunction& meshFunction, - const InFunction& function, - const RealType& time, - const RealType& outFunctionMultiplicator, - const RealType& inFunctionMultiplicator, - EntitisType entitiesType ) +evaluateEntities( OutMeshFunction& meshFunction, + const InFunction& function, + const RealType& time, + const RealType& outFunctionMultiplicator, + const RealType& inFunctionMultiplicator, + EntitiesType entitiesType ) { typedef typename MeshType::template MeshEntities< meshEntityDimensions > MeshEntityType; @@ -81,11 +84,10 @@ assign( OutMeshFunction& meshFunction, const EntityType& entity ) { typedef tnlFunctionAdapter< MeshType, InFunction > FunctionAdapter; - ( * userData.meshFunction )( entity ) = - FunctionAdapter::getValue( - *userData.function, - entity, - *userData.time ); + ( *userData.meshFunction )( entity ) = + *userData.outFunctionMultiplicator * ( *userData.meshFunction )( entity ) + + *userData.inFunctionMultiplicator * + FunctionAdapter::getValue( *userData.function, entity, *userData.time ); } }; @@ -160,6 +162,76 @@ assign( OutMeshFunction& meshFunction, } +template< typename OutMeshFunction, + typename Function, + typename Operator > +void +tnlMeshFunctionEvaluator< OutMeshFunction, tnlOperatorFunction< Operator, Function> >:: +evaluateEntities( OutMeshFunction& meshFunction, + const OperatorFunctionType& operatorFunction, + const RealType& time, + const RealType& outFunctionMultiplicator, + const RealType& inFunctionMultiplicator ) +{ + typedef typename MeshType::template MeshEntities< meshEntityDimensions > MeshEntityType; + + class AssignEntitiesProcessor + { + template< typename EntityType > + __cuda_callable__ + static inline void processEntity( const MeshType& mesh, + TraverserUserData& userData, + const EntityType& entity ) + { + typedef tnlFunctionAdapter< MeshType, InFunction > FunctionAdapter; + ( *userData.meshFunction )( entity ) = + *userData.outFunctionMultiplicator * ( *userData.meshFunction )( entity ) + + *userData.inFunctionMultiplicator * + FunctionAdapter::getValue( *userData.function, entity, *userData.time ); + } + }; + + if( std::is_same< MeshDeviceType, tnlHost >::value ) + { + TraverserUserData userData( &function, &time, &meshFunction, &outFunctionMultiplicator, &inFunctionMultiplicator ); + tnlTraverser< MeshType, MeshEntityType > meshTraverser; + meshTraverser.template processInterirorEntities< TraverserUserData, AssignEntitiesProcessor > + ( meshFunction.getMesh(), + userData ); + + } + if( std::is_same< MeshDeviceType, tnlCuda >::value ) + { + OutMeshFunction* kernelMeshFunction = tnlCuda::passToDevice( meshFunction ); + Function* kernelFunction = tnlCuda::passToDevice( *operatorFunction.function ); + Operator* kernelOperator = tnlCuda::passToDevice( *operatorFunction.operator_ ); + OperatorFunctionType auxOperatorFunction( *kernelOperator, *kernelFunction ); + OperatorFunctionType* kernelOperatorFunction = tnlCuda::passToDevice( auxOperatorFunction ); + RealType* kernelTime = tnlCuda::passToDevice( time ); + RealType* kernelOutFunctionMultiplicator = tnlCuda::passToDevice( outFunctionMultiplicator ); + RealType* kernelInFunctionMultiplicator = tnlCuda::passToDevice( inFunctionMultiplicator ); + + TraverserUserData userData( kernelOperatorFunction, kernelTime, kernelMeshFunction, kernelOutFunctionMultiplicator, kernelInFunctionMultiplicator ); + checkCudaDevice; + tnlTraverser< MeshType, MeshEntityType > meshTraverser; + meshTraverser.template processInteriorEntities< TraverserUserData, AssignEntitiesProcessor > + ( meshFunction.getMesh(), + userData ); + + + checkCudaDevice; + tnlCuda::freeFromDevice( kernelMeshFunction ); + tnlCuda::freeFromDevice( kernelFunction ); + tnlCuda::freeFromDevice( kernelOperator ); + tnlCuda::freeFromDevice( kernelOperatorFunction ); + tnlCuda::freeFromDevice( kernelTime ); + tnlCuda::freeFromDevice( kernelOutFunctionMultiplicator ); + tnlCuda::freeFromDevice( kernelInFunctionMultiplicator ); + + checkCudaDevice; + } +} + #endif /* TNLMESHFUNCTIONEVALUATOR_IMPL_H */ diff --git a/src/functions/tnlMeshFunction_impl.h b/src/functions/tnlMeshFunction_impl.h index 5f659d1650..678518dcc9 100644 --- a/src/functions/tnlMeshFunction_impl.h +++ b/src/functions/tnlMeshFunction_impl.h @@ -218,18 +218,7 @@ tnlMeshFunction< Mesh, MeshEntityDimensions, Real >& tnlMeshFunction< Mesh, MeshEntityDimensions, Real >:: operator = ( const Function& f ) { - if( std::is_void< typename Function::DeviceType >::value || - std::is_same< typename Function::DeviceType, DeviceType >::value ) - { - tnlFunctionEvaluator< ThisType, Function > evaluator; - evaluator.assignment( f, *this ); - return *this; - } - if( Function::getFunctionType() == MeshFunction ) - { - tnlMeshFunctionEvaluator< ThisType, Function >::assign( f, *this ); - return *this; - } + tnlMeshFunctionEvaluator< ThisType, Function >::evaluateAllEntities( *this, f ); return *this; } @@ -241,18 +230,7 @@ tnlMeshFunction< Mesh, MeshEntityDimensions, Real >& tnlMeshFunction< Mesh, MeshEntityDimensions, Real >:: operator += ( const Function& f ) { - if( std::is_void< typename Function::DeviceType >::value || - std::is_same< typename Function::DeviceType, DeviceType >::value ) - { - tnlFunctionEvaluator< ThisType, Function > evaluator; - evaluator.assignment( f, *this, 1.0, 1.0 ); - return *this; - } - if( Function::getFunctionType() == MeshFunction ) - { - tnlMeshFunctionEvaluator< ThisType, Function >::assign( f, *this ); - return *this; - } + tnlMeshFunctionEvaluator< ThisType, Function >::evaluateAllEntities( *this, f, 1.0, 1.0 ); return *this; } @@ -264,23 +242,10 @@ tnlMeshFunction< Mesh, MeshEntityDimensions, Real >& tnlMeshFunction< Mesh, MeshEntityDimensions, Real >:: operator -= ( const Function& f ) { - if( std::is_void< typename Function::DeviceType >::value || - std::is_same< typename Function::DeviceType, DeviceType >::value ) - { - tnlFunctionEvaluator< ThisType, Function > evaluator; - evaluator.assignment( f, *this ); - return *this; - } - if( Function::getFunctionType() == MeshFunction ) - { - tnlMeshFunctionEvaluator< ThisType, Function >::assign( f, *this, 1.0, -1.0 ); - return *this; - } + tnlMeshFunctionEvaluator< ThisType, Function >::evaluateAllEntities( *this, f, 1.0, -1.0 ); return *this; } - - template< typename Mesh, int MeshEntityDimensions, typename Real > diff --git a/src/functions/tnlOperatorFunction.h b/src/functions/tnlOperatorFunction.h index 19baae5cff..0caae80687 100644 --- a/src/functions/tnlOperatorFunction.h +++ b/src/functions/tnlOperatorFunction.h @@ -21,29 +21,25 @@ #include <type_traits> #include <core/tnlCuda.h> +/*** + * This class evaluates given operator on given function. The image is defined only + * for the interior mesh entities. + */ template< typename Operator, - typename Function, - typename BoundaryConditions = Operator > + typename Function > class tnlOperatorFunction -{ - //static_assert( std::is_same< typename Operator::MeshType, typename MeshFunction::MeshType >::value, - // "Operator and MeshFunction have different mesh type." ); - static_assert( std::is_same< typename Operator::MeshType, typename BoundaryConditions::MeshType >::value, - "Operator and BoundaryConditions have different mesh type." ); - +{ public: typedef Operator OperatorType; typedef Function FunctionType; - typedef BoundaryConditions BoundaryConditionsType; typedef typename OperatorType::RealType RealType; typedef typename OperatorType::DeviceType DeviceType; typedef typename OperatorType::IndexType IndexType; tnlOperatorFunction( const OperatorType& operator_, - const FunctionType& function, - const BoundaryConditionsType& boundaryConditions ); + const FunctionType& function ); template< typename MeshEntity > __cuda_callable__ @@ -53,14 +49,9 @@ class tnlOperatorFunction protected: - const Operator& operator_; + const Operator* operator_; - const FunctionType& function; - - const BoundaryConditions& boundaryConditions; - + const FunctionType* function; }; - - #endif /* TNLOPERATORFUNCTION_H */ diff --git a/src/functions/tnlOperatorFunction_impl.h b/src/functions/tnlOperatorFunction_impl.h index 978e4055c9..12ed6badee 100644 --- a/src/functions/tnlOperatorFunction_impl.h +++ b/src/functions/tnlOperatorFunction_impl.h @@ -24,11 +24,9 @@ template< typename Operator, tnlOperatorFunction< Operator, Function >:: tnlOperatorFunction( const OperatorType& operator_, - const MeshFunctionType& function, - const BoundaryConditionsType& boundaryConditions ) -: operator_( operator_ ), - function( function ), - boundaryConditions( boundaryConditions ) + const FunctionType& function ) +: operator_( &operator_ ), + function( &function ) { } @@ -41,12 +39,11 @@ tnlOperatorFunction< Operator, Function >::operator()( const MeshEntity& meshEntity, const RealType& time ) { - static_assert( MeshEntity::entityDimensions == Function::getMeshEntityDimensions(), "Wrong mesh entity dimensions." ); - if( meshEntity.isBoundaryEntity() ) - return boundaryConditions.getValue( meshEntity, function.getData(), time ); - return operator_.getValue( meshEntity, function.getData(), time ); + //static_assert( MeshEntity::entityDimensions == Operator::getMeshEntityDimensions(), "Wrong mesh entity dimensions." ); + tnlAssert( ! meshEntity.isBoundaryEntity(), + cerr << "Operator functions are defined only for interior mesh entities. Entity = " << meshEntity ); + return operator_->getValue( meshEntity, function->getData(), time ); } - #endif /* TNLOPERATORFUNCTION_IMPL_H */ diff --git a/src/operators/tnlOperatorEnumerator.h b/src/operators/tnlOperatorEnumerator.h deleted file mode 100644 index 050a23b72d..0000000000 --- a/src/operators/tnlOperatorEnumerator.h +++ /dev/null @@ -1,175 +0,0 @@ -/*************************************************************************** - tnlOperatorEnumerator.h - description - ------------------- - begin : Mar 8, 2015 - copyright : (C) 2015 by Tomas Oberhuber - email : tomas.oberhuber@fjfi.cvut.cz - ***************************************************************************/ - -/*************************************************************************** - * * - * This program is free software; you can redistribute it and/or modify * - * it under the terms of the GNU General Public License as published by * - * the Free Software Foundation; either version 2 of the License, or * - * (at your option) any later version. * - * * - ***************************************************************************/ -#ifndef SRC_OPERATORS_TNLOPERATORENUMERATOR_H_ -#define SRC_OPERATORS_TNLOPERATORENUMERATOR_H_ - -//#include <_operators/tnlOperatorAdapter.h> - -template< typename Operator, - typename DofVector > -class tnlOperatorEnumeratorTraverserUserData -{ - public: - - typedef typename DofVector::RealType RealType; - - const RealType *time; - - const Operator* _operator; - - DofVector *u; - - const RealType* _operatorCoefficient; - - const RealType* dofVectorCoefficient; - - tnlOperatorEnumeratorTraverserUserData( const RealType& time, - const Operator& _operator, - DofVector& u, - const RealType& _operatorCoefficient, - const RealType& dofVectorCoefficient ) - : time( &time ), - _operator( &_operator ), - u( &u ), - _operatorCoefficient( &_operatorCoefficient ), - dofVectorCoefficient( &dofVectorCoefficient ) - {}; -}; - - -template< typename Mesh, - typename Operator, - typename DofVector > -class tnlOperatorEnumerator -{ - public: - typedef Mesh MeshType; - typedef typename DofVector::RealType RealType; - typedef typename DofVector::DeviceType DeviceType; - typedef typename DofVector::IndexType IndexType; - typedef tnlOperatorEnumeratorTraverserUserData< Operator, - DofVector > TraverserUserData; - - template< int EntityDimensions > - void enumerate( const MeshType& mesh, - const Operator& _operator, - DofVector& u, - const RealType& _operatorCoefficient = 1.0, - const RealType& dofVectorCoefficient = 0.0, - const RealType& time = 0.0 ) const; - - - class TraverserEntitiesProcessor - { - public: - - template< int EntityDimensions > -#ifdef HAVE_CUDA - __host__ __device__ -#endif - static void processEntity( const MeshType& mesh, - TraverserUserData& userData, - const IndexType index ) - { - //typedef tnlOperatorAdapter< MeshType, Operator > OperatorAdapter; - ( *userData.u )[ index ] = - ( *userData.dofVectorCoefficient ) * ( *userData.u )[ index ] + - ( *userData._operatorCoefficient ) * userData._operator ->getValue( mesh, - index, - *userData.time ); - } - - }; - -}; - -template< int Dimensions, - typename Real, - typename Device, - typename Index, - typename Operator, - typename DofVector > -class tnlOperatorEnumerator< tnlGrid< Dimensions, Real, Device, Index >, - Operator, - DofVector > -{ - public: - - typedef tnlGrid< Dimensions, Real, Device, Index > MeshType; - typedef typename MeshType::RealType RealType; - typedef typename MeshType::DeviceType DeviceType; - typedef typename MeshType::IndexType IndexType; - typedef typename MeshType::CoordinatesType CoordinatesType; - typedef tnlOperatorEnumeratorTraverserUserData< Operator, - DofVector > TraverserUserData; - - template< int EntityDimensions > - void enumerate( const MeshType& mesh, - const Operator& _operator, - DofVector& u, - const RealType& _operatorCoefficient = 1.0, - const RealType& dofVectorCoefficient = 0.0, - const RealType& time = 0.0 ) const; - - class TraverserEntitiesProcessor - { - public: - - typedef typename MeshType::VertexType VertexType; - -#ifdef HAVE_CUDA - __host__ __device__ -#endif - static void processCell( const MeshType& mesh, - TraverserUserData& userData, - const IndexType index, - const CoordinatesType& coordinates ) - { - //printf( "Enumerator::processCell mesh =%p \n", &mesh ); - //typedef tnlOperatorAdapter< MeshType, Operator > OperatorAdapter; - ( *userData.u )[ index ] = - ( *userData.dofVectorCoefficient ) * ( *userData.u )[ index ] + - ( *userData._operatorCoefficient ) * userData._operator->getValue( mesh, - index, - coordinates, - *userData.time ); - - } - -#ifdef HAVE_CUDA - __host__ __device__ -#endif - static void processFace( const MeshType& mesh, - TraverserUserData& userData, - const IndexType index, - const CoordinatesType& coordinates ) - { - //typedef tnlOperatorAdapter< MeshType, Operator > OperatorAdapter; - ( *userData.u )[ index ] = - ( *userData.dofVectorCoefficient ) * ( *userData.u )[ index ] + - ( *userData._operatorCoefficient ) * userData._operator->getValue( mesh, - index, - coordinates, - *userData.time ); - } - }; - -}; - -#include <operators/tnlOperatorEnumerator_impl.h> - -#endif /* SRC_OPERATORS_TNLOPERATORENUMERATOR_H_ */ diff --git a/src/operators/tnlOperatorEnumerator_impl.h b/src/operators/tnlOperatorEnumerator_impl.h deleted file mode 100644 index ffe3f21a7a..0000000000 --- a/src/operators/tnlOperatorEnumerator_impl.h +++ /dev/null @@ -1,141 +0,0 @@ -/*************************************************************************** - tnlOperatorEnumerator_impl.h - description - ------------------- - begin : Mar 8, 2015 - copyright : (C) 2015 by Tomas Oberhuber - email : tomas.oberhuber@fjfi.cvut.cz - ***************************************************************************/ - -/*************************************************************************** - * * - * This program is free software; you can redistribute it and/or modify * - * it under the terms of the GNU General Public License as published by * - * the Free Software Foundation; either version 2 of the License, or * - * (at your option) any later version. * - * * - ***************************************************************************/ -#ifndef SRC_OPERATORS_TNLOPERATORENUMERATOR_IMPL_H_ -#define SRC_OPERATORS_TNLOPERATORENUMERATOR_IMPL_H_ - -#include <operators/tnlOperatorEnumerator.h> -#include <mesh/tnlTraverser_Grid1D.h> -#include <mesh/tnlTraverser_Grid2D.h> -#include <mesh/tnlTraverser_Grid3D.h> - -template< typename Mesh, - typename Operator, - typename DofVector > - template< int EntityDimensions > -void -tnlOperatorEnumerator< Mesh, Operator, DofVector >:: -enumerate( const MeshType& mesh, - const Operator& _operator, - DofVector& u, - const RealType& _operatorCoefficient, - const RealType& dofVectorCoefficient, - const RealType& time ) const - -{ - if( DeviceType::DeviceType == tnlHostDevice ) - { - TraverserUserData userData( time, _operator, u, _operatorCoefficient, dofVectorCoefficient ); - tnlTraverser< MeshType, EntityDimensions > meshTraverser; - meshTraverser.template processBoundaryEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - meshTraverser.template processInteriorEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - - } - if( DeviceType::DeviceType == tnlCudaDevice ) - { - RealType* kernelTime = tnlCuda::passToDevice( time ); - Operator* kernelOperator = tnlCuda::passToDevice( _operator ); - DofVector* kernelU = tnlCuda::passToDevice( u ); - RealType* kernelOperatorCoefficient = tnlCuda::passToDevice( _operatorCoefficient ); - RealType* kernelDofVectorCoefficient = tnlCuda::passToDevice( dofVectorCoefficient ); - TraverserUserData userData( *kernelTime, *kernelOperator, *kernelU, *kernelOperatorCoefficient, *kernelDofVectorCoefficient ); - checkCudaDevice; - tnlTraverser< MeshType, EntityDimensions > meshTraverser; - meshTraverser.template processBoundaryEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - meshTraverser.template processInteriorEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - - checkCudaDevice; - tnlCuda::freeFromDevice( kernelTime ); - tnlCuda::freeFromDevice( kernelOperator ); - tnlCuda::freeFromDevice( kernelU ); - tnlCuda::freeFromDevice( kernelOperatorCoefficient ); - tnlCuda::freeFromDevice( kernelDofVectorCoefficient ); - checkCudaDevice; - } -} - -template< int Dimensions, - typename Real, - typename Device, - typename Index, - typename Operator, - typename DofVector > - template< int EntityDimensions > -void -tnlOperatorEnumerator< tnlGrid< Dimensions, Real, Device, Index >, Operator, DofVector >:: -enumerate( const tnlGrid< Dimensions, Real, Device, Index >& mesh, - const Operator& _operator, - DofVector& u, - const RealType& _operatorCoefficient, - const RealType& dofVectorCoefficient, - const RealType& time ) const -{ - if( DeviceType::DeviceType == tnlHostDevice ) - { - TraverserUserData userData( time, _operator, u, _operatorCoefficient, dofVectorCoefficient ); - tnlTraverser< MeshType, EntityDimensions > meshTraverser; - meshTraverser.template processBoundaryEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - meshTraverser.template processInteriorEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - - } - if( DeviceType::DeviceType == tnlCudaDevice ) - { - RealType* kernelTime = tnlCuda::passToDevice( time ); - Operator* kernelOperator = tnlCuda::passToDevice( _operator ); - DofVector* kernelU = tnlCuda::passToDevice( u ); - RealType* kernelOperatorCoefficient = tnlCuda::passToDevice( _operatorCoefficient ); - RealType* kernelDofVectorCoefficient = tnlCuda::passToDevice( dofVectorCoefficient ); - TraverserUserData userData( *kernelTime, *kernelOperator, *kernelU, *kernelOperatorCoefficient, *kernelDofVectorCoefficient ); - checkCudaDevice; - tnlTraverser< MeshType, EntityDimensions > meshTraverser; - meshTraverser.template processBoundaryEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - meshTraverser.template processInteriorEntities< TraverserUserData, - TraverserEntitiesProcessor > - ( mesh, - userData ); - - checkCudaDevice; - tnlCuda::freeFromDevice( kernelTime ); - tnlCuda::freeFromDevice( kernelOperator ); - tnlCuda::freeFromDevice( kernelU ); - tnlCuda::freeFromDevice( kernelOperatorCoefficient ); - tnlCuda::freeFromDevice( kernelDofVectorCoefficient ); - checkCudaDevice; - } -} - -#endif /* SRC_OPERATORS_TNLOPERATORENUMERATOR_IMPL_H_ */ diff --git a/tests/unit-tests/tnlApproximationError_impl.h b/tests/unit-tests/tnlApproximationError_impl.h index 1317e024d0..dc04ed6738 100644 --- a/tests/unit-tests/tnlApproximationError_impl.h +++ b/tests/unit-tests/tnlApproximationError_impl.h @@ -46,12 +46,12 @@ getError( const ExactOperator& exactOperator, { typedef tnlMeshFunction< MeshType > MeshFunction; typedef tnlDirichletBoundaryConditions< MeshType, tnlConstantFunction< MeshType::meshDimensions > > BoundaryConditions; - typedef tnlOperatorFunction< ApproximateOperator, Function, BoundaryConditions > OperatorFunction; + typedef tnlOperatorFunction< ApproximateOperator, Function > OperatorFunction; typedef tnlExactOperatorFunction< ExactOperator, Function > ExactOperatorFunction; tnlMeshFunction< MeshType > exactU( mesh ), u( mesh ), v( mesh ); BoundaryConditions boundaryConditions; - OperatorFunction operatorFunction( approximateOperator, function, boundaryConditions ); + OperatorFunction operatorFunction( approximateOperator, function ); ExactOperatorFunction exactOperatorFunction( exactOperator, function ); exactU = exactOperatorFunction; -- GitLab