Commit 6fd14031 authored by Jakub Klinkovský's avatar Jakub Klinkovský
Browse files

Optimized CUDA traversers for cells and vertices 

Entity basis and orientation is passed only to faces and edges.
parent 27a5c94b

src/TNL/Meshes/GridDetails/GridTraverser.h

+24 −20
Original line number Diff line number Diff line
@@ -50,8 +50,6 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::Host, Index > > 
         const GridPointer& gridPointer,

         const CoordinatesType begin,

         const CoordinatesType end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

};
@@ -81,8 +79,6 @@ class GridTraverser< Meshes::Grid< 1, Real, Devices::Cuda, Index > > 
         const GridPointer& gridPointer,

         const CoordinatesType& begin,

         const CoordinatesType& end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

};
@@ -108,15 +104,17 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::Host, Index > > 
         typename UserData,

         bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary = 1,

         int YOrthogonalBoundary = 1 >

         int YOrthogonalBoundary = 1,

         typename... GridEntityParameters >

      static void

      processEntities(

         const GridPointer& gridPointer,

         const CoordinatesType begin,

         const CoordinatesType end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

         SharedPointer< UserData, DeviceType >& userData,

         // gridEntityParameters are passed to GridEntity's constructor

         // (i.e. orientation and basis for faces)

         const GridEntityParameters&... gridEntityParameters );

};



/****
@@ -141,15 +139,17 @@ class GridTraverser< Meshes::Grid< 2, Real, Devices::Cuda, Index > > 
         typename UserData,

         bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary = 1,

         int YOrthogonalBoundary = 1  >

         int YOrthogonalBoundary = 1,

         typename... GridEntityParameters >

      static void

      processEntities(

         const GridPointer& gridPointer,

         const CoordinatesType& begin,

         const CoordinatesType& end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

         SharedPointer< UserData, DeviceType >& userData,

         // gridEntityParameters are passed to GridEntity's constructor

         // (i.e. orientation and basis for faces)

         const GridEntityParameters&... gridEntityParameters );

};



/****
@@ -175,15 +175,17 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::Host, Index > > 
         bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary = 1,

         int YOrthogonalBoundary = 1,

         int ZOrthogonalBoundary = 1 >

         int ZOrthogonalBoundary = 1,

         typename... GridEntityParameters >

      static void

      processEntities(

         const GridPointer& gridPointer,

         const CoordinatesType begin,

         const CoordinatesType end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

         SharedPointer< UserData, DeviceType >& userData,

         // gridEntityParameters are passed to GridEntity's constructor

         // (i.e. orientation and basis for faces and edges)

         const GridEntityParameters&... gridEntityParameters );

};



/****
@@ -209,15 +211,17 @@ class GridTraverser< Meshes::Grid< 3, Real, Devices::Cuda, Index > > 
         bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary = 1,

         int YOrthogonalBoundary = 1,

         int ZOrthogonalBoundary = 1 >

         int ZOrthogonalBoundary = 1,

         typename... GridEntityParameters >

      static void

      processEntities(

         const GridPointer& gridPointer,

         const CoordinatesType& begin,

         const CoordinatesType& end,

         const CoordinatesType& entityOrientation,

         const CoordinatesType& entityBasis,

         SharedPointer< UserData, DeviceType >& userData );

         SharedPointer< UserData, DeviceType >& userData,

         // gridEntityParameters are passed to GridEntity's constructor

         // (i.e. orientation and basis for faces and edges)

         const GridEntityParameters&... gridEntityParameters );

};



} // namespace Meshes

src/TNL/Meshes/GridDetails/GridTraverser_impl.h

+62 −113
Original line number Diff line number Diff line
@@ -10,32 +10,9 @@ 


#pragma once



#include <TNL/UniquePointer.h>





namespace TNL {

namespace Meshes {



template< typename CoordinatesType >

struct TraverserKernelData

{

   CoordinatesType begin;

   CoordinatesType end;

   CoordinatesType entityOrientation;

   CoordinatesType entityBasis;



   TraverserKernelData( CoordinatesType begin,

                        CoordinatesType end,

                        CoordinatesType entityOrientation,

                        CoordinatesType entityBasis )

   : begin( begin ),

     end( end ),

     entityOrientation( entityOrientation ),

     entityBasis( entityBasis )

   {}

};





/****

 * 1D traverser, host

 */
@@ -52,18 +29,12 @@ processEntities( 
   const GridPointer& gridPointer,

   const CoordinatesType begin,

   const CoordinatesType end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

{

   GridEntity entity( *gridPointer );

   entity.setOrientation( entityOrientation );

   entity.setBasis( entityBasis );

   if( processOnlyBoundaryEntities )

   {

      GridEntity entity( *gridPointer );

      entity.setOrientation( entityOrientation );

      entity.setBasis( entityBasis );



      entity.getCoordinates() = begin;

      entity.refresh();
@@ -88,8 +59,6 @@ processEntities( 
#endif

      {

         GridEntity entity( *gridPointer );

         entity.setOrientation( entityOrientation );

         entity.setBasis( entityBasis );

#ifdef HAVE_OPENMP

#pragma omp for 

#endif
@@ -117,7 +86,8 @@ __global__ void 
GridTraverser1D(

   const Meshes::Grid< 1, Real, Devices::Cuda, Index >* grid,

   UserData* userData,

   const TraverserKernelData< typename GridEntity::CoordinatesType >* kernelData,

   const typename GridEntity::CoordinatesType begin,

   const typename GridEntity::CoordinatesType end,

   const Index gridIdx )

{

   typedef Real RealType;
@@ -125,10 +95,10 @@ GridTraverser1D( 
   typedef Meshes::Grid< 1, Real, Devices::Cuda, Index > GridType;

   typename GridType::CoordinatesType coordinates;

 

   coordinates.x() = kernelData->begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   if( coordinates.x() <= kernelData->end.x() )

   coordinates.x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   if( coordinates <= end )

   {   

      GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );

      GridEntity entity( *grid, coordinates );

      entity.refresh();

      EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );

   }
@@ -143,7 +113,8 @@ __global__ void 
GridBoundaryTraverser1D(

   const Meshes::Grid< 1, Real, Devices::Cuda, Index >* grid,

   UserData* userData,

   const TraverserKernelData< typename GridEntity::CoordinatesType >* kernelData )

   const typename GridEntity::CoordinatesType begin,

   const typename GridEntity::CoordinatesType end )

{

   typedef Real RealType;

   typedef Index IndexType;
@@ -152,15 +123,15 @@ GridBoundaryTraverser1D( 
 

   if( threadIdx.x == 0 )

   {

      coordinates.x() = kernelData->begin.x();

      GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );

      coordinates.x() = begin.x();

      GridEntity entity( *grid, coordinates );

      entity.refresh();

      EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );

   }

   if( threadIdx.x == 1 )

   {

      coordinates.x() = kernelData->end.x();

      GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );

      coordinates.x() = end.x();

      GridEntity entity( *grid, coordinates );

      entity.refresh();

      EntitiesProcessor::processEntity( entity.getMesh(), *userData, entity );

   }
@@ -181,14 +152,9 @@ processEntities( 
   const GridPointer& gridPointer,

   const CoordinatesType& begin,

   const CoordinatesType& end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

{

#ifdef HAVE_CUDA

   UniquePointer< TraverserKernelData< CoordinatesType >, Devices::Cuda >

      kernelData( begin, end, entityOrientation, entityBasis );



   Devices::Cuda::synchronizeDevice();

   if( processOnlyBoundaryEntities )

   {
@@ -198,7 +164,8 @@ processEntities( 
            <<< cudaBlocks, cudaBlockSize >>>

            ( &gridPointer.template getData< Devices::Cuda >(),

              &userDataPointer.template modifyData< Devices::Cuda >(),

              &kernelData.template getData< Devices::Cuda >() );

              begin,

              end );

   }

   else

   {
@@ -212,7 +179,8 @@ processEntities( 
            <<< cudaBlocks, cudaBlockSize >>>

            ( &gridPointer.template getData< Devices::Cuda >(),

              &userDataPointer.template modifyData< Devices::Cuda >(),

              &kernelData.template getData< Devices::Cuda >(),

              begin,

              end,

              gridXIdx );

   }

   cudaThreadSynchronize();
@@ -232,22 +200,20 @@ template< typename Real, 
      typename UserData,

      bool processOnlyBoundaryEntities,

      int XOrthogonalBoundary,

      int YOrthogonalBoundary >

      int YOrthogonalBoundary,

      typename... GridEntityParameters >

void

GridTraverser< Meshes::Grid< 2, Real, Devices::Host, Index > >::

processEntities(

   const GridPointer& gridPointer,

   const CoordinatesType begin,

   const CoordinatesType end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

   SharedPointer< UserData, DeviceType >& userDataPointer,

   const GridEntityParameters&... gridEntityParameters )

{

   if( processOnlyBoundaryEntities )

   {

      GridEntity entity( *gridPointer );

      entity.setOrientation( entityOrientation );

      entity.setBasis( entityBasis );

      GridEntity entity( *gridPointer, begin, gridEntityParameters... );

      

      if( YOrthogonalBoundary )

         for( entity.getCoordinates().x() = begin.x();
@@ -292,9 +258,7 @@ processEntities( 
#pragma omp parallel firstprivate( begin, end ) if( Devices::Host::isOMPEnabled() )

#endif

      {

         GridEntity entity( *gridPointer );

         entity.setOrientation( entityOrientation );

         entity.setBasis( entityBasis );

         GridEntity entity( *gridPointer, begin, gridEntityParameters... );

#ifdef HAVE_OPENMP

#pragma omp for 

#endif
@@ -319,33 +283,27 @@ template< typename Real, 
          typename GridEntity,

          typename UserData,

          typename EntitiesProcessor,

          bool processOnlyBoundaryEntities >

          bool processOnlyBoundaryEntities,

          typename... GridEntityParameters >

__global__ void 

GridTraverser2D(

   const Meshes::Grid< 2, Real, Devices::Cuda, Index >* grid,

   UserData* userData,

   //const TraverserKernelData< typename GridEntity::CoordinatesType >* kernelData,

   const typename GridEntity::CoordinatesType begin,

   const typename GridEntity::CoordinatesType end,

   const typename GridEntity::CoordinatesType entityOrientation,

   const typename GridEntity::CoordinatesType entityBasis,

   const Index gridXIdx,

   const Index gridYIdx )

   const Index gridYIdx,

   const GridEntityParameters... gridEntityParameters )

{

   typedef Meshes::Grid< 2, Real, Devices::Cuda, Index > GridType;

   typename GridType::CoordinatesType coordinates;



   //coordinates.x() = kernelData->begin.x() + ( gridXIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   //coordinates.y() = kernelData->begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;  

   

   coordinates.x() = begin.x() + ( gridXIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   coordinates.y() = begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;  





   if( coordinates.x() <= end.x() &&

       coordinates.y() <= end.y() )

   if( coordinates <= end )

   {

      GridEntity entity( *grid, coordinates, entityOrientation, entityBasis );

      GridEntity entity( *grid, coordinates, gridEntityParameters... );

      entity.refresh();

      if( ! processOnlyBoundaryEntities || entity.isBoundaryEntity() )

      {
@@ -366,21 +324,18 @@ template< typename Real, 
      typename UserData,

      bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary,

         int YOrthogonalBoundary >

         int YOrthogonalBoundary,

      typename... GridEntityParameters >

void

GridTraverser< Meshes::Grid< 2, Real, Devices::Cuda, Index > >::

processEntities(

   const GridPointer& gridPointer,

   const CoordinatesType& begin,

   const CoordinatesType& end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

   SharedPointer< UserData, DeviceType >& userDataPointer,

   const GridEntityParameters&... gridEntityParameters )

{

#ifdef HAVE_CUDA   

   //UniquePointer< TraverserKernelData< CoordinatesType >, Devices::Cuda >

   //   kernelData( begin, end, entityOrientation, entityBasis );



   dim3 cudaBlockSize( 16, 16 );

   dim3 cudaBlocks;

   cudaBlocks.x = Devices::Cuda::getNumberOfBlocks( end.x() - begin.x() + 1, cudaBlockSize.x );
@@ -391,14 +346,15 @@ processEntities( 
   Devices::Cuda::synchronizeDevice();

   for( IndexType gridYIdx = 0; gridYIdx < cudaYGrids; gridYIdx ++ )

      for( IndexType gridXIdx = 0; gridXIdx < cudaXGrids; gridXIdx ++ )

         GridTraverser2D< Real, Index, GridEntity, UserData, EntitiesProcessor, processOnlyBoundaryEntities >

         GridTraverser2D< Real, Index, GridEntity, UserData, EntitiesProcessor, processOnlyBoundaryEntities, GridEntityParameters... >

            <<< cudaBlocks, cudaBlockSize >>>

            ( &gridPointer.template getData< Devices::Cuda >(),

              &userDataPointer.template modifyData< Devices::Cuda >(),

              //&kernelData.template getData< Devices::Cuda >(),

              begin, end, entityOrientation, entityBasis,

              begin,

              end,

              gridXIdx,

              gridYIdx );

              gridYIdx,

              gridEntityParameters... );

 

   cudaThreadSynchronize();

   checkCudaDevice;
@@ -417,22 +373,20 @@ template< typename Real, 
      bool processOnlyBoundaryEntities,

      int XOrthogonalBoundary,

      int YOrthogonalBoundary,

      int ZOrthogonalBoundary >

      int ZOrthogonalBoundary,

      typename... GridEntityParameters >

void

GridTraverser< Meshes::Grid< 3, Real, Devices::Host, Index > >::

processEntities(

   const GridPointer& gridPointer,

   const CoordinatesType begin,

   const CoordinatesType end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

   SharedPointer< UserData, DeviceType >& userDataPointer,

   const GridEntityParameters&... gridEntityParameters )

{

   if( processOnlyBoundaryEntities )

   {

      GridEntity entity( *gridPointer );

      entity.setOrientation( entityOrientation );

      entity.setBasis( entityBasis );

      GridEntity entity( *gridPointer, begin, gridEntityParameters... );

      

      if( ZOrthogonalBoundary )

         for( entity.getCoordinates().y() = begin.y();
@@ -501,9 +455,7 @@ processEntities( 
#pragma omp parallel firstprivate( begin, end ) if( Devices::Host::isOMPEnabled() )

#endif

      {

         GridEntity entity( *gridPointer );

         entity.setOrientation( entityOrientation );

         entity.setBasis( entityBasis );         

         GridEntity entity( *gridPointer, begin, gridEntityParameters... );

#ifdef HAVE_OPENMP

#pragma omp for

#endif
@@ -530,30 +482,29 @@ template< typename Real, 
          typename GridEntity,

          typename UserData,

          typename EntitiesProcessor,

          bool processOnlyBoundaryEntities >

          bool processOnlyBoundaryEntities,

          typename... GridEntityParameters >

__global__ void

GridTraverser3D(

   const Meshes::Grid< 3, Real, Devices::Cuda, Index >* grid,

   UserData* userData,

   const TraverserKernelData< typename GridEntity::CoordinatesType >* kernelData,

   const typename GridEntity::CoordinatesType begin,

   const typename GridEntity::CoordinatesType end,

   const Index gridXIdx,

   const Index gridYIdx,

   const Index gridZIdx )

   const Index gridZIdx,

   const GridEntityParameters... gridEntityParameters )

{

   typedef Meshes::Grid< 3, Real, Devices::Cuda, Index > GridType;

   typename GridType::CoordinatesType coordinates;



   coordinates.x() = kernelData->begin.x() + ( gridXIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   coordinates.y() = kernelData->begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;

   coordinates.z() = kernelData->begin.z() + ( gridZIdx * Devices::Cuda::getMaxGridSize() + blockIdx.z ) * blockDim.z + threadIdx.z;

 

   

   coordinates.x() = begin.x() + ( gridXIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;

   coordinates.y() = begin.y() + ( gridYIdx * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y;

   coordinates.z() = begin.z() + ( gridZIdx * Devices::Cuda::getMaxGridSize() + blockIdx.z ) * blockDim.z + threadIdx.z;



   if( coordinates.x() <= kernelData->end.x() &&

       coordinates.y() <= kernelData->end.y() &&

       coordinates.z() <= kernelData->end.z() )

   if( coordinates <= end )

   {

      GridEntity entity( *grid, coordinates, kernelData->entityOrientation, kernelData->entityBasis );

      GridEntity entity( *grid, coordinates, gridEntityParameters... );

      entity.refresh();

      if( ! processOnlyBoundaryEntities || entity.isBoundaryEntity() )

      {
@@ -575,21 +526,18 @@ template< typename Real, 
      bool processOnlyBoundaryEntities,

         int XOrthogonalBoundary,

         int YOrthogonalBoundary,

         int ZOrthogonalBoundary >

         int ZOrthogonalBoundary,

      typename... GridEntityParameters >

void

GridTraverser< Meshes::Grid< 3, Real, Devices::Cuda, Index > >::

processEntities(

   const GridPointer& gridPointer,

   const CoordinatesType& begin,

   const CoordinatesType& end,

   const CoordinatesType& entityOrientation,

   const CoordinatesType& entityBasis,

   SharedPointer< UserData, DeviceType >& userDataPointer )

   SharedPointer< UserData, DeviceType >& userDataPointer,

   const GridEntityParameters&... gridEntityParameters )

{

#ifdef HAVE_CUDA   

   UniquePointer< TraverserKernelData< CoordinatesType >, Devices::Cuda >

      kernelData( begin, end, entityOrientation, entityBasis );

      

   dim3 cudaBlockSize( 8, 8, 8 );

   dim3 cudaBlocks;

   cudaBlocks.x = Devices::Cuda::getNumberOfBlocks( end.x() - begin.x() + 1, cudaBlockSize.x );
@@ -603,14 +551,16 @@ processEntities( 
   for( IndexType gridZIdx = 0; gridZIdx < cudaZGrids; gridZIdx ++ )

      for( IndexType gridYIdx = 0; gridYIdx < cudaYGrids; gridYIdx ++ )

         for( IndexType gridXIdx = 0; gridXIdx < cudaXGrids; gridXIdx ++ )

            GridTraverser3D< Real, Index, GridEntity, UserData, EntitiesProcessor, processOnlyBoundaryEntities >

            GridTraverser3D< Real, Index, GridEntity, UserData, EntitiesProcessor, processOnlyBoundaryEntities, GridEntityParameters... >

               <<< cudaBlocks, cudaBlockSize >>>

               ( &gridPointer.template getData< Devices::Cuda >(),

                 &userDataPointer.template modifyData< Devices::Cuda >(),

                 &kernelData.template getData< Devices::Cuda >(),

                 begin,

                 end,

                 gridXIdx,

                 gridYIdx,

                 gridZIdx );

                 gridZIdx,

                 gridEntityParameters... );



   cudaThreadSynchronize();

   checkCudaDevice;
@@ -619,4 +569,3 @@ processEntities( 


} // namespace Meshes

} // namespace TNL

src/TNL/Meshes/GridDetails/Traverser_Grid1D_impl.h

+0 −12
Original line number Diff line number Diff line
@@ -38,8 +38,6 @@ processBoundaryEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -63,8 +61,6 @@ processInteriorEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 1 ),

      gridPointer->getDimensions() - CoordinatesType( 2 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -89,8 +85,6 @@ processAllEntities( 
      gridPointer,

      CoordinatesType( 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -117,8 +111,6 @@ processBoundaryEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0 ),

      gridPointer->getDimensions(),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -142,8 +134,6 @@ processInteriorEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 1 ),

      gridPointer->getDimensions() - CoordinatesType( 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -168,8 +158,6 @@ processAllEntities( 
      gridPointer,

      CoordinatesType( 0 ),

      gridPointer->getDimensions(),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}

src/TNL/Meshes/GridDetails/Traverser_Grid2D_impl.h

+18 −30
Original line number Diff line number Diff line
@@ -39,8 +39,6 @@ processBoundaryEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -64,8 +62,6 @@ processInteriorEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 1, 1 ),

      gridPointer->getDimensions() - CoordinatesType( 2, 2 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -89,8 +85,6 @@ processAllEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -113,21 +107,21 @@ processBoundaryEntities( const GridPointer& gridPointer, 
    */

   static_assert( GridEntity::entityDimensions == 1, "The entity has wrong dimensions." );

 

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, true, 1, 0 >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, true, 1, 0, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 0, 1 ),

      userDataPointer,

      CoordinatesType( 1, 0 ),

      CoordinatesType( 0, 1 ),

      userDataPointer );

      CoordinatesType( 0, 1 ) );



   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, true, 0, 1 >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, true, 0, 1, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 0 ),

      userDataPointer,

      CoordinatesType( 0, 1 ),

      CoordinatesType( 1, 0 ),

      userDataPointer );

      CoordinatesType( 1, 0 ) );

}



template< typename Real,
@@ -146,21 +140,21 @@ processInteriorEntities( const GridPointer& gridPointer, 
    */

   static_assert( GridEntity::entityDimensions == 1, "The entity has wrong dimensions." );

 

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 1, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 1 ),

      userDataPointer,

      CoordinatesType( 1, 0 ),

      CoordinatesType( 0, 1 ),

      userDataPointer );

      CoordinatesType( 0, 1 ) );



   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 0, 1 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 1 ),

      userDataPointer,

      CoordinatesType( 0, 1 ),

      CoordinatesType( 1, 0 ),

      userDataPointer );

      CoordinatesType( 1, 0 ) );

}



template< typename Real,
@@ -179,21 +173,21 @@ processAllEntities( const GridPointer& gridPointer, 
    */

   static_assert( GridEntity::entityDimensions == 1, "The entity has wrong dimensions." );

 

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 0, 1 ),

      userDataPointer,

      CoordinatesType( 1, 0 ),

      CoordinatesType( 0, 1 ),

      userDataPointer );

      CoordinatesType( 0, 1 ) );



   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false >(

   GridTraverser< GridType >::template processEntities< GridEntity, EntitiesProcessor, UserData, false, CoordinatesType, CoordinatesType >(

      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 0 ),

      userDataPointer,

      CoordinatesType( 0, 1 ),

      CoordinatesType( 1, 0 ),

      userDataPointer );

      CoordinatesType( 1, 0 ) );

}



template< typename Real,
@@ -216,8 +210,6 @@ processBoundaryEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions(),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -241,8 +233,6 @@ processInteriorEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 1, 1 ),

      gridPointer->getDimensions() - CoordinatesType( 1, 1 ),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}
@@ -266,8 +256,6 @@ processAllEntities( const GridPointer& gridPointer, 
      gridPointer,

      CoordinatesType( 0, 0 ),

      gridPointer->getDimensions(),

      CoordinatesType(),

      CoordinatesType(),

      userDataPointer );

}

src/TNL/Meshes/GridDetails/Traverser_Grid3D_impl.h

+54 −66

File changed.

Preview size limit exceeded, changes collapsed.