From 16cf2aa74001635b5a360da03ccd026cc0841234 Mon Sep 17 00:00:00 2001
From: Tomas Oberhuber <tomas.oberhuber@fjfi.cvut.cz>
Date: Sun, 8 Mar 2015 18:21:33 +0100
Subject: [PATCH] Adding operator enumerator.
---
TODO | 2 +
src/functions/tnlFunctionEnumerator.h | 1 +
src/mesh/tnlTraverser_Grid2D_impl.h | 2 -
src/operators/CMakeLists.txt | 4 +-
src/operators/tnlOperatorEnumerator.h | 175 +++++++++++++++++++++
src/operators/tnlOperatorEnumerator_impl.h | 141 +++++++++++++++++
6 files changed, 322 insertions(+), 3 deletions(-)
create mode 100644 src/operators/tnlOperatorEnumerator.h
create mode 100644 src/operators/tnlOperatorEnumerator_impl.h
diff --git a/TODO b/TODO
index 410049b00d..c8ed457124 100644
--- a/TODO
+++ b/TODO
@@ -1,3 +1,5 @@
+TODO: v tnlMeshResolver se provadi preklad pro vsechny mozne sablonove parametry => prorezat
+
TODO: napsat FunctionDiscretizer pro jednotne rozhrani RightHandSide
TODO: doplnit mesh travelsals pro jine mesh entity nez cell
diff --git a/src/functions/tnlFunctionEnumerator.h b/src/functions/tnlFunctionEnumerator.h
index 9f238b8660..ad96e49ece 100644
--- a/src/functions/tnlFunctionEnumerator.h
+++ b/src/functions/tnlFunctionEnumerator.h
@@ -140,6 +140,7 @@ class tnlFunctionEnumerator< tnlGrid< Dimensions, Real, Device, Index >,
const IndexType index,
const CoordinatesType& coordinates )
{
+ //printf( "Enumerator::processCell mesh =%p \n", &mesh );
typedef tnlFunctionAdapter< MeshType, Function > FunctionAdapter;
( *userData.u )[ index ] =
( *userData.dofVectorCoefficient ) * ( *userData.u )[ index ] +
diff --git a/src/mesh/tnlTraverser_Grid2D_impl.h b/src/mesh/tnlTraverser_Grid2D_impl.h
index a40a20c965..6000b1af7a 100644
--- a/src/mesh/tnlTraverser_Grid2D_impl.h
+++ b/src/mesh/tnlTraverser_Grid2D_impl.h
@@ -352,8 +352,6 @@ processBoundaryEntities( const GridType& grid,
GridType* kernelGrid = tnlCuda::passToDevice( grid );
UserData* kernelUserData = tnlCuda::passToDevice( userData );
- printf( "Number of cells = %d \n", grid.getNumberOfCells() );
-
dim3 cudaBlockSize( 16, 16 );
dim3 cudaBlocks;
cudaBlocks.x = tnlCuda::getNumberOfBlocks( grid.getDimensions().x(), cudaBlockSize.x );
diff --git a/src/operators/CMakeLists.txt b/src/operators/CMakeLists.txt
index 8d6c37b489..6ddc8f02ab 100755
--- a/src/operators/CMakeLists.txt
+++ b/src/operators/CMakeLists.txt
@@ -12,7 +12,9 @@ SET( headers tnlFiniteDifferences.h
tnlNeumannBoundaryConditions_impl.h
tnlAnalyticNeumannBoundaryConditions.h
tnlAnalyticNeumannBoundaryConditions_impl.h
- tnlExactOperatorEvaluator.h )
+ tnlExactOperatorEvaluator.h
+ tnlOperatorEnumerator.h
+ tnlOperatorEnumerator_impl.h )
SET( CURRENT_DIR ${CMAKE_SOURCE_DIR}/src/operators )
diff --git a/src/operators/tnlOperatorEnumerator.h b/src/operators/tnlOperatorEnumerator.h
new file mode 100644
index 0000000000..050a23b72d
--- /dev/null
+++ b/src/operators/tnlOperatorEnumerator.h
@@ -0,0 +1,175 @@
+/***************************************************************************
+ 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
new file mode 100644
index 0000000000..ffe3f21a7a
--- /dev/null
+++ b/src/operators/tnlOperatorEnumerator_impl.h
@@ -0,0 +1,141 @@
+/***************************************************************************
+ 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_ */
--
GitLab