From 1ace5365d2dc74120e49ed2adaa9b0ffa76bf4e8 Mon Sep 17 00:00:00 2001
From: =?UTF-8?q?Tom=C3=A1=C5=A1=20Oberhuber?= <oberhuber.tomas@gmail.com>
Date: Sat, 5 Jan 2019 08:30:26 +0100
Subject: [PATCH] Added synchrounous/asynchronous modes for grid traversers.
---
src/TNL/Meshes/GridDetails/GridTraverser.h | 29 ++++++++++++----
.../Meshes/GridDetails/GridTraverser_1D.hpp | 18 +++++++---
.../Meshes/GridDetails/GridTraverser_2D.hpp | 20 +++++++----
.../Meshes/GridDetails/GridTraverser_3D.hpp | 13 ++++---
.../GridDetails/Traverser_Grid1D_impl.h | 30 ++++++++++------
.../GridDetails/Traverser_Grid2D_impl.h | 18 ++++++++++
.../GridDetails/Traverser_Grid3D_impl.h | 34 ++++++++++++++++++-
7 files changed, 129 insertions(+), 33 deletions(-)
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser.h b/src/TNL/Meshes/GridDetails/GridTraverser.h
index 881367d3f1..fb6b34da12 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser.h
+++ b/src/TNL/Meshes/GridDetails/GridTraverser.h
@@ -25,6 +25,8 @@ class GridTraverser
{
};
+enum GridTraverserMode { synchronousMode, asynchronousMode };
+
/****
* 1D grid, Devices::Host
*/
@@ -52,6 +54,7 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::Host, Index > >
const CoordinatesType begin,
const CoordinatesType end,
UserData& userData,
+ GridTraverserMode mode = synchronousMode,
const int& stream = 0 );
};
@@ -82,6 +85,7 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::Cuda, Index > >
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode = synchronousMode,
const int& stream = 0 );
};
@@ -112,6 +116,7 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::MIC, Index > >
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode = synchronousMode,
const int& stream = 0 );
};
@@ -148,7 +153,9 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::Host, Index > >
const CoordinatesType end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces)
@@ -186,7 +193,9 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::Cuda, Index > >
const CoordinatesType& end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces)
@@ -224,7 +233,9 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::MIC, Index > >
const CoordinatesType& end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces)
@@ -263,7 +274,9 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::Host, Index > >
const CoordinatesType end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces and edges)
@@ -302,7 +315,9 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::Cuda, Index > >
const CoordinatesType& end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces and edges)
@@ -341,7 +356,9 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::MIC, Index > >
const CoordinatesType& end,
UserData& userData,
// FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
-// const int& stream = 0,
+ //GridTraverserMode mode = synchronousMode,
+ GridTraverserMode mode,
+ // const int& stream = 0,
const int& stream,
// gridEntityParameters are passed to GridEntity's constructor
// (i.e. orientation and basis for faces and edges)
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
index 90148f8e8d..505f9c3d74 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_1D.hpp
@@ -41,6 +41,7 @@ processEntities(
const CoordinatesType begin,
const CoordinatesType end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream )
{
GridEntity entity( *gridPointer );
@@ -177,13 +178,14 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream )
{
#ifdef HAVE_CUDA
auto& pool = CudaStreamPool::getInstance();
const cudaStream_t& s = pool.getStream( stream );
- Devices::Cuda::synchronizeDevice();
+ //Devices::Cuda::synchronizeDevice();
if( processOnlyBoundaryEntities )
{
dim3 cudaBlockSize( 2 );
@@ -209,15 +211,20 @@ processEntities(
userData,
begin,
end,
- gridXIdx );
+ gridXIdx );*/
}
- // only launches into the stream 0 are synchronized
- /*if( stream == 0 )
+#ifdef NDEBUG
+ if( mode == synchronousMode )
{
cudaStreamSynchronize( s );
TNL_CHECK_CUDA_DEVICE;
- }*/
+ }
+#else
+ cudaStreamSynchronize( s );
+ TNL_CHECK_CUDA_DEVICE;
+#endif
+
#else
throw Exceptions::CudaSupportMissing();
#endif
@@ -241,6 +248,7 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream )
{
std::cout << "Not Implemented yet Grid Traverser <1, Real, Device::MIC>" << std::endl;
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
index 84e496017f..50b30c0190 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_2D.hpp
@@ -43,6 +43,7 @@ processEntities(
const CoordinatesType begin,
const CoordinatesType end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
@@ -402,6 +403,7 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
@@ -534,13 +536,18 @@ processEntities(
gridEntityParameters... );
}
- // only launches into the stream 0 are synchronized
- if( stream == 0 )
- {
- cudaStreamSynchronize( s );
- TNL_CHECK_CUDA_DEVICE;
- }
+#ifdef NDEBUG
+ if( mode == synchronousMode )
+ {
+ cudaStreamSynchronize( s );
+ TNL_CHECK_CUDA_DEVICE;
}
+#else
+ cudaStreamSynchronize( s );
+ TNL_CHECK_CUDA_DEVICE;
+#endif
+ }
+
#else
throw Exceptions::CudaSupportMissing();
#endif
@@ -567,6 +574,7 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
diff --git a/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp b/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
index d63b81f46a..9259da9bf4 100644
--- a/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
+++ b/src/TNL/Meshes/GridDetails/GridTraverser_3D.hpp
@@ -42,6 +42,7 @@ processEntities(
const CoordinatesType begin,
const CoordinatesType end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
@@ -324,6 +325,7 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
@@ -390,7 +392,7 @@ processEntities(
( &gridPointer.template getData< Devices::Cuda >(),
userData,
begin.x(),
- end.x(),
+ end.x(),
begin.z() + 1,
end.z() - 1,
begin.y(),
@@ -401,7 +403,7 @@ processEntities(
( &gridPointer.template getData< Devices::Cuda >(),
userData,
begin.x(),
- end.x(),
+ end.x(),
begin.z() + 1,
end.z() - 1,
end.y(),
@@ -417,7 +419,7 @@ processEntities(
( &gridPointer.template getData< Devices::Cuda >(),
userData,
begin.y() + 1,
- end.y() - 1,
+ end.y() - 1,
begin.z() + 1,
end.z() - 1,
begin.x(),
@@ -428,7 +430,7 @@ processEntities(
( &gridPointer.template getData< Devices::Cuda >(),
userData,
begin.y() + 1,
- end.y() - 1,
+ end.y() - 1,
begin.z() + 1,
end.z() - 1,
end.x(),
@@ -440,7 +442,7 @@ processEntities(
cudaStreamSynchronize( s3 );
cudaStreamSynchronize( s4 );
cudaStreamSynchronize( s5 );
- cudaStreamSynchronize( s6 );
+ cudaStreamSynchronize( s6 );
TNL_CHECK_CUDA_DEVICE;
}
else
@@ -506,6 +508,7 @@ processEntities(
const CoordinatesType& begin,
const CoordinatesType& end,
UserData& userData,
+ GridTraverserMode mode,
const int& stream,
const GridEntityParameters&... gridEntityParameters )
{
diff --git a/src/TNL/Meshes/GridDetails/Traverser_Grid1D_impl.h b/src/TNL/Meshes/GridDetails/Traverser_Grid1D_impl.h
index 99ea85876a..5669f6e83e 100644
--- a/src/TNL/Meshes/GridDetails/Traverser_Grid1D_impl.h
+++ b/src/TNL/Meshes/GridDetails/Traverser_Grid1D_impl.h
@@ -43,7 +43,8 @@ processBoundaryEntities( const GridPointer& gridPointer,
gridPointer,
CoordinatesType( 0 ),
gridPointer->getDimensions() - CoordinatesType( 1 ),
- userData );
+ userData,
+ asynchronousMode );
}
else //Distributed
{
@@ -54,7 +55,8 @@ processBoundaryEntities( const GridPointer& gridPointer,
gridPointer,
CoordinatesType( 0 ) + distributedGrid->getLowerOverlap(),
CoordinatesType( 0 ) + distributedGrid->getLowerOverlap(),
- userData );
+ userData,
+ asynchronousMode );
}
if( neighbors[ Meshes::DistributedMeshes::Right ] == -1 )
@@ -63,7 +65,8 @@ processBoundaryEntities( const GridPointer& gridPointer,
gridPointer,
gridPointer->getDimensions() - CoordinatesType( 1 ) - distributedGrid->getUpperOverlap(),
gridPointer->getDimensions() - CoordinatesType( 1 ) - distributedGrid->getUpperOverlap(),
- userData );
+ userData,
+ asynchronousMode );
}
}
@@ -92,7 +95,8 @@ processInteriorEntities( const GridPointer& gridPointer,
gridPointer,
CoordinatesType( 1 ),
gridPointer->getDimensions() - CoordinatesType( 2 ),
- userData );
+ userData,
+ asynchronousMode );
}
else //Distributed
{
@@ -117,7 +121,8 @@ processInteriorEntities( const GridPointer& gridPointer,
gridPointer,
begin,
end,
- userData );
+ userData,
+ asynchronousMode );
}
}
@@ -146,7 +151,8 @@ processAllEntities(
gridPointer,
CoordinatesType( 0 ),
gridPointer->getDimensions() - CoordinatesType( 1 ),
- userData );
+ userData,
+ asynchronousMode );
}
else //Distributed
{
@@ -157,7 +163,8 @@ processAllEntities(
gridPointer,
begin,
end,
- userData );
+ userData,
+ asynchronousMode );
}
}
@@ -185,7 +192,8 @@ processBoundaryEntities( const GridPointer& gridPointer,
gridPointer,
CoordinatesType( 0 ),
gridPointer->getDimensions(),
- userData );
+ userData,
+ asynchronousMode );
}
template< typename Real,
@@ -208,7 +216,8 @@ processInteriorEntities( const GridPointer& gridPointer,
gridPointer,
CoordinatesType( 1 ),
gridPointer->getDimensions() - CoordinatesType( 1 ),
- userData );
+ userData,
+ asynchronousMode );
}
template< typename Real,
@@ -232,7 +241,8 @@ processAllEntities(
gridPointer,
CoordinatesType( 0 ),
gridPointer->getDimensions(),
- userData );
+ userData,
+ asynchronousMode );
}
} // namespace Meshes
diff --git a/src/TNL/Meshes/GridDetails/Traverser_Grid2D_impl.h b/src/TNL/Meshes/GridDetails/Traverser_Grid2D_impl.h
index 23d93d7e0b..4d87b18ba0 100644
--- a/src/TNL/Meshes/GridDetails/Traverser_Grid2D_impl.h
+++ b/src/TNL/Meshes/GridDetails/Traverser_Grid2D_impl.h
@@ -42,6 +42,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
else //Distributed
@@ -57,6 +58,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
begin,
CoordinatesType( begin.x(), end.y() ),
userData,
+ asynchronousMode,
0 );
}
@@ -67,6 +69,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( end.x(), begin.y() ),
end,
userData,
+ asynchronousMode,
0 );
}
@@ -78,6 +81,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
begin,
CoordinatesType( end.x(), begin.y() ),
userData,
+ asynchronousMode,
0 );
}
@@ -88,6 +92,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( begin.x(), end.y() ),
end,
userData,
+ asynchronousMode,
0 );
}
}
@@ -117,6 +122,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 1 ),
gridPointer->getDimensions() - CoordinatesType( 2, 2 ),
userData,
+ asynchronousMode,
0 );
}
else // distributed
@@ -142,6 +148,7 @@ processInteriorEntities( const GridPointer& gridPointer,
begin,
end,
userData,
+ asynchronousMode,
0);
}
}
@@ -170,6 +177,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
else
@@ -183,6 +191,7 @@ processAllEntities( const GridPointer& gridPointer,
begin,
end,
userData,
+ asynchronousMode,
0);
}
}
@@ -211,6 +220,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0 ),
CoordinatesType( 0, 1 ) );
@@ -220,6 +230,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 1 ),
CoordinatesType( 1, 0 ) );
@@ -246,6 +257,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0 ),
CoordinatesType( 0, 1 ) );
@@ -255,6 +267,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 1 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 1 ),
CoordinatesType( 1, 0 ) );
@@ -281,6 +294,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0 ),
CoordinatesType( 0, 1 ) );
@@ -290,6 +304,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 1 ),
CoordinatesType( 1, 0 ) );
@@ -316,6 +331,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions(),
userData,
+ asynchronousMode,
0 );
}
@@ -340,6 +356,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 1 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
@@ -364,6 +381,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0 ),
gridPointer->getDimensions(),
userData,
+ asynchronousMode,
0 );
}
diff --git a/src/TNL/Meshes/GridDetails/Traverser_Grid3D_impl.h b/src/TNL/Meshes/GridDetails/Traverser_Grid3D_impl.h
index 3c9fffd813..f4575dfec8 100644
--- a/src/TNL/Meshes/GridDetails/Traverser_Grid3D_impl.h
+++ b/src/TNL/Meshes/GridDetails/Traverser_Grid3D_impl.h
@@ -44,6 +44,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
else // distributed
@@ -59,6 +60,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
begin,
CoordinatesType( begin.x(), end.y(), end.z() ),
userData,
+ asynchronousMode,
0 );
}
@@ -69,6 +71,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( end.x() , begin.y(), begin.z() ),
end,
userData,
+ asynchronousMode,
0 );
}
@@ -79,6 +82,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
begin,
CoordinatesType( end.x(), begin.y(), end.z() ),
userData,
+ asynchronousMode,
0 );
}
@@ -89,6 +93,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( begin.x(), end.y(), begin.z() ),
end,
userData,
+ asynchronousMode,
0 );
}
@@ -99,6 +104,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
begin,
CoordinatesType( end.x(), end.y(), begin.z() ),
userData,
+ asynchronousMode,
0 );
}
@@ -109,6 +115,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( begin.x(), begin.y(), end.z() ),
end,
userData,
+ asynchronousMode,
0 );
}
}
@@ -138,6 +145,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 1, 1 ),
gridPointer->getDimensions() - CoordinatesType( 2, 2, 2 ),
userData,
+ asynchronousMode,
0 );
}
else
@@ -169,7 +177,8 @@ processInteriorEntities( const GridPointer& gridPointer,
begin,
end,
userData,
- 0);
+ asynchronousMode,
+ 0 );
}
}
@@ -197,6 +206,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
else
@@ -209,6 +219,7 @@ processAllEntities( const GridPointer& gridPointer,
begin,
end,
userData,
+ asynchronousMode,
0 );
}
}
@@ -237,6 +248,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1, 1 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 1, 0, 0 ),
CoordinatesType( 0, 1, 1 ) );
@@ -246,6 +258,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 0, 1, 0 ),
CoordinatesType( 1, 0, 1 ) );
@@ -255,6 +268,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 0 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 0, 1 ),
CoordinatesType( 1, 1, 0 ) );
@@ -281,6 +295,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 1, 0, 0 ),
CoordinatesType( 0, 1, 1 ) );
@@ -290,6 +305,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 1, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 0, 1, 0 ),
CoordinatesType( 1, 0, 1 ) );
@@ -299,6 +315,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 1 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 0, 1 ),
CoordinatesType( 1, 1, 0 ) );
@@ -324,6 +341,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1, 1 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 1, 0, 0 ),
CoordinatesType( 0, 1, 1 ) );
@@ -333,6 +351,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 0, 1, 0 ),
CoordinatesType( 1, 0, 1 ) );
@@ -342,6 +361,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 0 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 0, 0, 1 ),
CoordinatesType( 1, 1, 0 ) );
@@ -371,6 +391,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0, 0 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 0, 1, 1 ),
CoordinatesType( 1, 0, 0 ) );
@@ -380,6 +401,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1, 0 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0, 1 ),
CoordinatesType( 0, 1, 0 ) );
@@ -389,6 +411,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 0, 1 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 1, 1, 0 ),
CoordinatesType( 0, 0, 1 ) );
@@ -415,6 +438,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 1, 1 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1, 1 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 0, 1, 1 ),
CoordinatesType( 1, 0, 0 ) );
@@ -424,6 +448,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 0, 1 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0, 1 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0, 1 ),
CoordinatesType( 0, 1, 0 ) );
@@ -433,6 +458,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 1, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 0 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 1, 1, 0 ),
CoordinatesType( 0, 0, 1 ) );
@@ -458,6 +484,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 1, 0, 0 ),
userData,
+ asynchronousMode,
2,
CoordinatesType( 0, 1, 1 ),
CoordinatesType( 1, 0, 0 ) );
@@ -467,6 +494,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 1, 0 ),
userData,
+ asynchronousMode,
1,
CoordinatesType( 1, 0, 1 ),
CoordinatesType( 0, 1, 0 ) );
@@ -476,6 +504,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions() - CoordinatesType( 0, 0, 1 ),
userData,
+ asynchronousMode,
0,
CoordinatesType( 1, 1, 0 ),
CoordinatesType( 0, 0, 1 ) );
@@ -505,6 +534,7 @@ processBoundaryEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions(),
userData,
+ asynchronousMode,
0 );
}
@@ -529,6 +559,7 @@ processInteriorEntities( const GridPointer& gridPointer,
CoordinatesType( 1, 1, 1 ),
gridPointer->getDimensions() - CoordinatesType( 1, 1, 1 ),
userData,
+ asynchronousMode,
0 );
}
@@ -553,6 +584,7 @@ processAllEntities( const GridPointer& gridPointer,
CoordinatesType( 0, 0, 0 ),
gridPointer->getDimensions(),
userData,
+ asynchronousMode,
0 );
}
--
GitLab