From af8983d7c28d5513795af0f636401dbc8c2899c9 Mon Sep 17 00:00:00 2001
From: Tomas Oberhuber <tomas.oberhuber@fjfi.cvut.cz>
Date: Sat, 28 Feb 2015 12:45:03 +0100
Subject: [PATCH] Adding face traversing to the grid traverser.
---
install | 2 +-
src/mesh/tnlTraverser_Grid2D_impl.h | 12 +++---
src/solvers/pde/tnlExplicitUpdater.h | 38 +++++++++++++++++
src/solvers/pde/tnlLinearSystemAssembler.h | 48 ++++++++++++++++++++++
4 files changed, 93 insertions(+), 7 deletions(-)
diff --git a/install b/install
index 217830e0b6..0c15ff9b1e 100755
--- a/install
+++ b/install
@@ -2,7 +2,7 @@
TARGET=TNL
INSTALL_PREFIX=${HOME}/local
-WITH_CUDA=yes
+WITH_CUDA=no
TEMPLATE_EXPLICIT_INSTANTIATION=yes
#VERBOSE="VERBOSE=1"
diff --git a/src/mesh/tnlTraverser_Grid2D_impl.h b/src/mesh/tnlTraverser_Grid2D_impl.h
index 60105b57ee..b765b265da 100644
--- a/src/mesh/tnlTraverser_Grid2D_impl.h
+++ b/src/mesh/tnlTraverser_Grid2D_impl.h
@@ -97,16 +97,16 @@ processBoundaryEntities( const GridType& grid,
for( coordinates.x() = 0; coordinates.x() < xSize; coordinates.x() ++ )
{
coordinates.y() = 0;
- EntitiesProcessor::processFace< 0, 1 >( grid, userData, grid.getFaceIndex< 0, 1 >( coordinates ), coordinates );
+ EntitiesProcessor::processFace( grid, userData, grid.getFaceIndex< 0, 1 >( coordinates ), coordinates );
coordinates.y() = ySize;
- EntitiesProcessor::processFace< 0, 1 >( grid, userData, grid.getFaceIndex< 0, 1 >( coordinates ), coordinates );
+ EntitiesProcessor::processFace( grid, userData, grid.getFaceIndex< 0, 1 >( coordinates ), coordinates );
}
for( coordinates.y() = 0; coordinates.y() < ySize; coordinates.y() ++ )
{
coordinates.x() = 0;
- EntitiesProcessor::processFace< 1, 0 >( grid, userData, grid.getFaceIndex< 1, 0 >( coordinates ), coordinates );
+ EntitiesProcessor::processFace( grid, userData, grid.getFaceIndex< 1, 0 >( coordinates ), coordinates );
coordinates.x() = ySize;
- EntitiesProcessor::processFace< 1, 0 >( grid, userData, grid.getFaceIndex< 1, 0 >( coordinates ), coordinates );
+ EntitiesProcessor::processFace( grid, userData, grid.getFaceIndex< 1, 0 >( coordinates ), coordinates );
}
}
@@ -134,13 +134,13 @@ processInteriorEntities( const GridType& grid,
for( coordinates.x() = 1; coordinates.x() < xSize; coordinates.x() ++ )
{
const IndexType index = grid.template getFaceIndex< 1, 0 >( coordinates );
- EntitiesProcessor::processFace< 1, 0 >( grid, userData, index, coordinates );
+ EntitiesProcessor::processFace( grid, userData, index, coordinates );
}
for( coordinates.y() = 1; coordinates.y() < ySize; coordinates.y() ++ )
for( coordinates.x() = 1; coordinates.x() < xSize - 1; coordinates.x() ++ )
{
const IndexType index = grid.template getFaceIndex< 0, 1 >( coordinates );
- EntitiesProcessor::processFace< 0, 1 >( grid, userData, index, coordinates );
+ EntitiesProcessor::processFace( grid, userData, index, coordinates );
}
}
diff --git a/src/solvers/pde/tnlExplicitUpdater.h b/src/solvers/pde/tnlExplicitUpdater.h
index 01432ace7f..e77536157a 100644
--- a/src/solvers/pde/tnlExplicitUpdater.h
+++ b/src/solvers/pde/tnlExplicitUpdater.h
@@ -190,6 +190,23 @@ class tnlExplicitUpdater< tnlGrid< Dimensions, Real, Device, Index >,
*userData.fu );
}
+#ifdef HAVE_CUDA
+ __host__ __device__
+#endif
+ static void processFace( const MeshType& mesh,
+ TraversalUserData& userData,
+ const IndexType index,
+ const CoordinatesType& coordinates )
+ {
+ userData.boundaryConditions->setBoundaryConditions( *userData.time,
+ mesh,
+ index,
+ coordinates,
+ *userData.u,
+ *userData.fu );
+ }
+
+
};
class TraversalInteriorEntitiesProcessor
@@ -220,6 +237,27 @@ class tnlExplicitUpdater< tnlGrid< Dimensions, Real, Device, Index >,
*userData.time );
}
+#ifdef HAVE_CUDA
+ __host__ __device__
+#endif
+ static void processFace( const MeshType& mesh,
+ TraversalUserData& userData,
+ const IndexType index,
+ const CoordinatesType& coordinates )
+ {
+ ( *userData.fu)[ index ] = userData.differentialOperator->getValue( mesh,
+ index,
+ coordinates,
+ *userData.u,
+ *userData.time );
+
+ typedef tnlFunctionAdapter< MeshType, RightHandSide > FunctionAdapter;
+ ( * userData.fu )[ index ] += FunctionAdapter::getValue( mesh,
+ *userData.rightHandSide,
+ index,
+ coordinates,
+ *userData.time );
+ }
};
};
diff --git a/src/solvers/pde/tnlLinearSystemAssembler.h b/src/solvers/pde/tnlLinearSystemAssembler.h
index 0898e1e293..01beb9065a 100644
--- a/src/solvers/pde/tnlLinearSystemAssembler.h
+++ b/src/solvers/pde/tnlLinearSystemAssembler.h
@@ -227,6 +227,26 @@ class tnlLinearSystemAssembler< tnlGrid< Dimensions, Real, Device, Index >,
matrixRow );
}
+#ifdef HAVE_CUDA
+ __host__ __device__
+#endif
+ static void processFace( const MeshType& mesh,
+ TraversalUserData& userData,
+ const IndexType index,
+ const CoordinatesType& coordinates )
+ {
+ //printf( "index = %d \n", index );
+ typename MatrixType::MatrixRow matrixRow = userData.matrix->getRow( index );
+ userData.boundaryConditions->updateLinearSystem( *userData.time,
+ mesh,
+ index,
+ coordinates,
+ *userData.u,
+ *userData.b,
+ matrixRow );
+ }
+
+
};
class TraversalInteriorEntitiesProcessor
@@ -262,6 +282,34 @@ class tnlLinearSystemAssembler< tnlGrid< Dimensions, Real, Device, Index >,
userData.matrix->addElementFast( index, index, 1.0, 1.0 );
}
+#ifdef HAVE_CUDA
+ __host__ __device__
+#endif
+ static void processFace( const MeshType& mesh,
+ TraversalUserData& userData,
+ const IndexType index,
+ const CoordinatesType& coordinates )
+ {
+ //printf( "index = %d \n", index );
+ typedef tnlFunctionAdapter< MeshType, RightHandSide > FunctionAdapter;
+ ( *userData.b )[ index ] = ( *userData.u )[ index ] +
+ ( *userData.tau ) * FunctionAdapter::getValue( mesh,
+ *userData.rightHandSide,
+ index,
+ coordinates,
+ *userData.time );
+
+ typename MatrixType::MatrixRow matrixRow = userData.matrix->getRow( index );
+ userData.differentialOperator->updateLinearSystem( *userData.time,
+ *userData.tau,
+ mesh,
+ index,
+ coordinates,
+ *userData.u,
+ *userData.b,
+ matrixRow );
+ userData.matrix->addElementFast( index, index, 1.0, 1.0 );
+ }
};
};
--
GitLab