diff --git a/install b/install index 217830e0b630bc3bb0a67552eec438508640725c..0c15ff9b1e3e6f64833becbfd12d2b004e53e2c2 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 60105b57ee83fd213e4f93f7a31426b78280b6ab..b765b265daa4089e535761df06a70dfb7fd85b00 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 01432ace7fba41d153cf9f8b0520a959018e54f7..e77536157aa3b2b58fc53cc9b2fde430c61222d1 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 0898e1e2938534beaf071b4297a65fbe10a27ba4..01beb9065a9f63afdcca069e4a74448459b35bd7 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 ); + } }; };