Commit 806d6e76 authored by Tomáš Oberhuber's avatar Tomáš Oberhuber
Browse files

Fixes after rebase - deleting useless files.

parent b2a22b31
Loading
Loading
Loading
Loading
+0 −114
Original line number Diff line number Diff line
// Copyright (c) 2004-2022 Tomáš Oberhuber et al.
//
// This file is part of TNL - Template Numerical Library (https://tnl-project.org/)
//
// SPDX-License-Identifier: MIT

#pragma once

#include <TNL/Meshes/Grid.h>
#include <TNL/Pointers/SharedPointer.h>

namespace TNL {
namespace Meshes {

/****
 * This is only a helper class for Traverser specializations for Grid.
 */
template< typename Grid >
class GridTraverser
{};

enum GridTraverserMode
{
   synchronousMode,
   asynchronousMode
};

enum GridTraverserMode { synchronousMode, asynchronousMode };


/****
 * 3D grid, Devices::Host
 */
template< typename Real, typename Index >
class GridTraverser< Meshes::Grid< 3, Real, Devices::Host, Index > >
{
   public:

      typedef Meshes::Grid< 3, Real, Devices::Host, Index > GridType;
      typedef Pointers::SharedPointer<  GridType > GridPointer;
      typedef Real RealType;
      typedef Devices::Host DeviceType;
      typedef Index IndexType;
      typedef typename GridType::CoordinatesType CoordinatesType;

      template<
         typename GridEntity,
         typename EntitiesProcessor,
         typename UserData,
         bool processOnlyBoundaryEntities,
         int XOrthogonalBoundary = 1,
         int YOrthogonalBoundary = 1,
         int ZOrthogonalBoundary = 1,
         typename... GridEntityParameters >
      static void
      processEntities(
         const GridPointer& gridPointer,
         const CoordinatesType begin,
         const CoordinatesType end,
         UserData& userData,
         // FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
         //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)
         const GridEntityParameters&... gridEntityParameters );
};

/****
 * 3D grid, Devices::Cuda
 */
template< typename Real, typename Index >
class GridTraverser< Meshes::Grid< 3, Real, Devices::Cuda, Index > >
{
   public:

      typedef Meshes::Grid< 3, Real, Devices::Cuda, Index > GridType;
      typedef Pointers::SharedPointer<  GridType > GridPointer;
      typedef Real RealType;
      typedef Devices::Cuda DeviceType;
      typedef Index IndexType;
      typedef typename GridType::CoordinatesType CoordinatesType;

      template<
         typename GridEntity,
         typename EntitiesProcessor,
         typename UserData,
         bool processOnlyBoundaryEntities,
         int XOrthogonalBoundary = 1,
         int YOrthogonalBoundary = 1,
         int ZOrthogonalBoundary = 1,
         typename... GridEntityParameters >
      static void
      processEntities(
         const GridPointer& gridPointer,
         const CoordinatesType& begin,
         const CoordinatesType& end,
         UserData& userData,
         // FIXME: hack around nvcc bug (error: default argument not at end of parameter list)
         //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)
         const GridEntityParameters&... gridEntityParameters );
};

}  // namespace Meshes
}  // namespace TNL

#include <TNL/Meshes/GridDetails/GridTraverser_3D.hpp>
+0 −197
Original line number Diff line number Diff line
// Copyright (c) 2004-2022 Tomáš Oberhuber et al.
//
// This file is part of TNL - Template Numerical Library (https://tnl-project.org/)
//
// SPDX-License-Identifier: MIT

// Implemented by: Tomas Oberhuber,
//                 Jakub Klinkovsky,
//                 Vit Hanousek

#pragma once

#include <TNL/Meshes/Grid.h>
#include <TNL/Pointers/SharedPointer.h>
#include <TNL/Cuda/StreamPool.h>
#include <TNL/Exceptions/CudaSupportMissing.h>
#include <TNL/Meshes/GridDetails/GridTraverser.h>
#include <TNL/Exceptions/NotImplementedError.h>

namespace TNL {
namespace Meshes {

/****
 * 1D traverser, host
 */
template< typename Real, typename Index >
template< typename GridEntity, typename EntitiesProcessor, typename UserData, bool processOnlyBoundaryEntities >
void
GridTraverser< Meshes::Grid< 1, Real, Devices::Host, Index > >::processEntities( const GridPointer& gridPointer,
                                                                                 const CoordinatesType& begin,
                                                                                 const CoordinatesType& end,
                                                                                 UserData& userData,
                                                                                 GridTraverserMode mode,
                                                                                 const int& stream )
{
   GridEntity entity( *gridPointer );
   if( processOnlyBoundaryEntities ) {
      GridEntity entity( *gridPointer );

      entity.getCoordinates() = begin;
      entity.refresh();
      EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
      entity.getCoordinates() = end - 1;
      entity.refresh();
      EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
   }
   else {
#ifdef HAVE_OPENMP
      if( Devices::Host::isOMPEnabled() && end.x() - begin.x() > 512 ) {
         #pragma omp parallel firstprivate( begin, end )
         {
            GridEntity entity( *gridPointer );
            #pragma omp for
            // TODO: g++ 5.5 crashes when coding this loop without auxiliary x as bellow
            for( IndexType x = begin.x(); x < end.x(); x++ ) {
               entity.getCoordinates().x() = x;
               entity.refresh();
               EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
            }
         }
      }
      else {
         GridEntity entity( *gridPointer );
         for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() < end.x(); entity.getCoordinates().x()++ ) {
            entity.refresh();
            EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
         }
      }
#else
      GridEntity entity( *gridPointer );
      for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() < end.x(); entity.getCoordinates().x()++ ) {
         entity.refresh();
         EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
      }
#endif
   }
}

/****
 * 1D traverser, CUDA
 */
#ifdef HAVE_CUDA
template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor >
__global__
void
GridTraverser1D( const Meshes::Grid< 1, Real, Devices::Cuda, Index >* grid,
                 UserData userData,
                 const typename GridEntity::CoordinatesType begin,
                 const typename GridEntity::CoordinatesType end,
                 const Index gridIdx )
{
   typedef Real RealType;
   typedef Index IndexType;
   typedef Meshes::Grid< 1, Real, Devices::Cuda, Index > GridType;
   typename GridType::CoordinatesType coordinates;

   coordinates.x() = begin.x() + ( gridIdx * Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x;
   if( coordinates < end ) {
      GridEntity entity( *grid, coordinates );
      entity.refresh();
      EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
   }
}

template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor >
__global__
void
GridBoundaryTraverser1D( const Meshes::Grid< 1, Real, Devices::Cuda, Index >* grid,
                         UserData userData,
                         const typename GridEntity::CoordinatesType begin,
                         const typename GridEntity::CoordinatesType end )
{
   typedef Real RealType;
   typedef Index IndexType;
   typedef Meshes::Grid< 1, Real, Devices::Cuda, Index > GridType;
   typename GridType::CoordinatesType coordinates;

   if( threadIdx.x == 0 ) {
      coordinates.x() = begin.x();
      GridEntity entity( *grid, coordinates );
      entity.refresh();
      EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
   }
   if( threadIdx.x == 1 ) {
      coordinates.x() = end.x() - 1;
      GridEntity entity( *grid, coordinates );
      entity.refresh();
      EntitiesProcessor::processEntity( entity.getMesh(), userData, entity );
   }
}

#endif

template< typename Real, typename Index >
template< typename GridEntity, typename EntitiesProcessor, typename UserData, bool processOnlyBoundaryEntities >
void
GridTraverser< Meshes::Grid< 1, Real, Devices::Cuda, Index > >::processEntities( const GridPointer& gridPointer,
                                                                                 const CoordinatesType& begin,
                                                                                 const CoordinatesType& end,
                                                                                 UserData& userData,
                                                                                 GridTraverserMode mode,
                                                                                 const int& stream )
{
#ifdef HAVE_CUDA
   auto& pool = Cuda::StreamPool::getInstance();
   const cudaStream_t& s = pool.getStream( stream );

   Pointers::synchronizeSmartPointersOnDevice< Devices::Cuda >();
   if( processOnlyBoundaryEntities ) {
      dim3 cudaBlockSize( 2 );
      dim3 cudaBlocks( 1 );
      GridBoundaryTraverser1D< Real, Index, GridEntity, UserData, EntitiesProcessor > <<< cudaBlocks,
         cudaBlockSize, 0, s >>>( &gridPointer.template getData< Devices::Cuda >(), userData, begin, end );
   }
   else {
      dim3 blockSize( 256 ), blocksCount, gridsCount;
      Cuda::setupThreads( blockSize, blocksCount, gridsCount, end.x() - begin.x() );
      dim3 gridIdx;
      for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) {
         dim3 gridSize;
         Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize );
         GridTraverser1D< Real, Index, GridEntity, UserData, EntitiesProcessor > <<< blocksCount, blockSize, 0,
            s >>>( &gridPointer.template getData< Devices::Cuda >(), userData, begin, end, gridIdx.x );
      }

      /*dim3 cudaBlockSize( 256 );
      dim3 cudaBlocks;
      cudaBlocks.x = Cuda::getNumberOfBlocks( end.x() - begin.x(), cudaBlockSize.x );
      const IndexType cudaXGrids = Cuda::getNumberOfGrids( cudaBlocks.x );

      for( IndexType gridXIdx = 0; gridXIdx < cudaXGrids; gridXIdx ++ )
         GridTraverser1D< Real, Index, GridEntity, UserData, EntitiesProcessor >
            <<< cudaBlocks, cudaBlockSize, 0, s >>>
            ( &gridPointer.template getData< Devices::Cuda >(),
              userData,
              begin,
              end,
              gridXIdx );*/
   }

   #ifdef NDEBUG
   if( mode == synchronousMode ) {
      cudaStreamSynchronize( s );
      TNL_CHECK_CUDA_DEVICE;
   }
   #else
   cudaStreamSynchronize( s );
   TNL_CHECK_CUDA_DEVICE;
   #endif

#else
   throw Exceptions::CudaSupportMissing();
#endif
}

}  // namespace Meshes
}  // namespace TNL
+0 −524

File deleted.

Preview size limit exceeded, changes collapsed.

+0 −463

File deleted.

Preview size limit exceeded, changes collapsed.

+0 −101
Original line number Diff line number Diff line
// Copyright (c) 2004-2022 Tomáš Oberhuber et al.
//
// This file is part of TNL - Template Numerical Library (https://tnl-project.org/)
//
// SPDX-License-Identifier: MIT

#pragma once

#include <TNL/Cuda/CudaCallable.h>
#include <TNL/Meshes/DimensionTag.h>
#include <TNL/Meshes/GridEntityConfig.h>
#include <TNL/Meshes/GridDetails/NeighborGridEntityGetter.h>

namespace TNL {
namespace Meshes {

template< typename GridEntity,
          int NeighborEntityDimension,
          typename GridEntityConfig,
          bool storage = GridEntityConfig::template neighborEntityStorage< GridEntity >( NeighborEntityDimension ) >
class NeighborGridEntityLayer : public NeighborGridEntityLayer< GridEntity, NeighborEntityDimension - 1, GridEntityConfig >
{
public:
   using BaseType = NeighborGridEntityLayer< GridEntity, NeighborEntityDimension - 1, GridEntityConfig >;
   using NeighborEntityGetterType = NeighborGridEntityGetter< GridEntity, NeighborEntityDimension >;

   using BaseType::getNeighborEntities;

   __cuda_callable__
   NeighborGridEntityLayer( const GridEntity& entity ) : BaseType( entity ), neighborEntities( entity ) {}

   __cuda_callable__
   const NeighborEntityGetterType&
   getNeighborEntities( const DimensionTag< NeighborEntityDimension >& tag ) const
   {
      return this->neighborEntities;
   }

   __cuda_callable__
   void
   refresh( const typename GridEntity::GridType& grid, const typename GridEntity::GridType::IndexType& entityIndex )
   {
      BaseType::refresh( grid, entityIndex );
      neighborEntities.refresh( grid, entityIndex );
   }

protected:
   NeighborEntityGetterType neighborEntities;
};

template< typename GridEntity, typename GridEntityConfig, bool storage >
class NeighborGridEntityLayer< GridEntity, 0, GridEntityConfig, storage >
{
public:
   using NeighborEntityGetterType = NeighborGridEntityGetter< GridEntity, 0 >;

   __cuda_callable__
   NeighborGridEntityLayer( const GridEntity& entity ) : neighborEntities( entity ) {}

   __cuda_callable__
   const NeighborEntityGetterType&
   getNeighborEntities( const DimensionTag< 0 >& tag ) const
   {
      return this->neighborEntities;
   }

   __cuda_callable__
   void
   refresh( const typename GridEntity::GridType& grid, const typename GridEntity::GridType::IndexType& entityIndex )
   {
      neighborEntities.refresh( grid, entityIndex );
   }

protected:
   NeighborEntityGetterType neighborEntities;
};

template< typename GridEntity, typename GridEntityConfig >
class NeighborGridEntitiesStorage
: public NeighborGridEntityLayer< GridEntity, GridEntity::getMeshDimension(), GridEntityConfig >
{
   using BaseType = NeighborGridEntityLayer< GridEntity, GridEntity::getMeshDimension(), GridEntityConfig >;

public:
   using BaseType::getNeighborEntities;
   using BaseType::refresh;

   __cuda_callable__
   NeighborGridEntitiesStorage( const GridEntity& entity ) : BaseType( entity ) {}

   template< int EntityDimension >
   __cuda_callable__
   const NeighborGridEntityGetter< GridEntity, EntityDimension >&
   getNeighborEntities() const
   {
      return BaseType::getNeighborEntities( DimensionTag< EntityDimension >() );
   }
};

}  // namespace Meshes
}  // namespace TNL
Loading