Loading src/Benchmarks/Traversers/GridTraverserBenchmarkHelper.h +18 −14 Original line number Diff line number Diff line Loading @@ -37,18 +37,19 @@ _GridTraverser1D( typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 1, Real, Devices::Cuda, Index > GridType; typename GridType::CoordinatesType coordinates; //typename GridType::CoordinatesType coordinates; GridEntity entity;//( *grid, ); //entity.getCoordinates().x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; coordinates.x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; if( coordinates <= end ) GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; //coordinates.x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; if( entity.getCoordinates() <= end ) { //entity.refresh(); entity.refresh(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( RealType ) 1.0; userData.data[ coordinates.x() ] += ( RealType ) 1.0; //( *userData.u )( entity ) += ( RealType ) 1.0; //userData.data[ entity.getIndex() ] += ( RealType ) 1.0; //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //EntitiesProcessor::processEntity( entity.getMesh(), userData, entity ); } } Loading Loading @@ -80,22 +81,25 @@ class GridTraverserBenchmarkHelper< Grid, Devices::Host > UserDataType& userData, std::size_t size ) { /*Meshes::GridTraverser< Grid >::template processEntities< Cell, WriteOneEntitiesProcessorType, WriteOneTraverserUserDataType, false >( /*Meshes::GridTraverser< Grid >::template processEntities< CellType, AddOneEntitiesProcessorType, UserDataType, false >( grid, CoordinatesType( 0 ), grid->getDimensions() - CoordinatesType( 1 ), userData );*/ userData ); */ const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); //MeshFunction* _u = &u.template modifyData< Device >(); /*SimpleCellType entity( *grid ); SimpleCellType entity( *grid ); for( IndexType x = begin.x(); x <= end.x(); x ++ ) { entity.getCoordinates().x() = x; entity.refresh(); AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); }*/ //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); } } }; Loading src/Benchmarks/Traversers/SimpleCell.h 0 → 100644 +95 −0 Original line number Diff line number Diff line /*************************************************************************** SimpleCell.h - description ------------------- begin : Jan 5, 2019 copyright : (C) 2019 by oberhuber email : tomas.oberhuber@fjfi.cvut.cz ***************************************************************************/ /* See Copyright Notice in tnl/Copyright */ // Implemented by: Tomas Oberhuber #pragma once #include <TNL/Devices/Cuda.h> #include <TNL/Meshes/Grid.h> namespace TNL { namespace Benchmarks { namespace Traversers { template< typename Grid > class SimpleCell{}; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 1, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 1; }; __cuda_callable__ SimpleCell( const GridType& grid ) : grid( grid ){}; __cuda_callable__ const GridType& getMesh() const { return this->grid;}; __cuda_callable__ CoordinatesType& getCoordinates() { return this->coordinates; }; __cuda_callable__ void refresh() {index = coordinates.x();}; __cuda_callable__ const IndexType& getIndex() const { return this->index; }; protected: const GridType& grid; CoordinatesType coordinates; IndexType index; }; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 2, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 2; }; }; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 3, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 3; }; }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL Loading
src/Benchmarks/Traversers/GridTraverserBenchmarkHelper.h +18 −14 Original line number Diff line number Diff line Loading @@ -37,18 +37,19 @@ _GridTraverser1D( typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 1, Real, Devices::Cuda, Index > GridType; typename GridType::CoordinatesType coordinates; //typename GridType::CoordinatesType coordinates; GridEntity entity;//( *grid, ); //entity.getCoordinates().x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; coordinates.x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; if( coordinates <= end ) GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; //coordinates.x() = begin.x() + ( gridIdx * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; if( entity.getCoordinates() <= end ) { //entity.refresh(); entity.refresh(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( RealType ) 1.0; userData.data[ coordinates.x() ] += ( RealType ) 1.0; //( *userData.u )( entity ) += ( RealType ) 1.0; //userData.data[ entity.getIndex() ] += ( RealType ) 1.0; //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //EntitiesProcessor::processEntity( entity.getMesh(), userData, entity ); } } Loading Loading @@ -80,22 +81,25 @@ class GridTraverserBenchmarkHelper< Grid, Devices::Host > UserDataType& userData, std::size_t size ) { /*Meshes::GridTraverser< Grid >::template processEntities< Cell, WriteOneEntitiesProcessorType, WriteOneTraverserUserDataType, false >( /*Meshes::GridTraverser< Grid >::template processEntities< CellType, AddOneEntitiesProcessorType, UserDataType, false >( grid, CoordinatesType( 0 ), grid->getDimensions() - CoordinatesType( 1 ), userData );*/ userData ); */ const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); //MeshFunction* _u = &u.template modifyData< Device >(); /*SimpleCellType entity( *grid ); SimpleCellType entity( *grid ); for( IndexType x = begin.x(); x <= end.x(); x ++ ) { entity.getCoordinates().x() = x; entity.refresh(); AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); }*/ //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); } } }; Loading
src/Benchmarks/Traversers/SimpleCell.h 0 → 100644 +95 −0 Original line number Diff line number Diff line /*************************************************************************** SimpleCell.h - description ------------------- begin : Jan 5, 2019 copyright : (C) 2019 by oberhuber email : tomas.oberhuber@fjfi.cvut.cz ***************************************************************************/ /* See Copyright Notice in tnl/Copyright */ // Implemented by: Tomas Oberhuber #pragma once #include <TNL/Devices/Cuda.h> #include <TNL/Meshes/Grid.h> namespace TNL { namespace Benchmarks { namespace Traversers { template< typename Grid > class SimpleCell{}; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 1, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 1; }; __cuda_callable__ SimpleCell( const GridType& grid ) : grid( grid ){}; __cuda_callable__ const GridType& getMesh() const { return this->grid;}; __cuda_callable__ CoordinatesType& getCoordinates() { return this->coordinates; }; __cuda_callable__ void refresh() {index = coordinates.x();}; __cuda_callable__ const IndexType& getIndex() const { return this->index; }; protected: const GridType& grid; CoordinatesType coordinates; IndexType index; }; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 2, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 2; }; }; template< typename Real, typename Device, typename Index > class SimpleCell< Meshes::Grid< 3, Real, Device, Index > > { public: using GridType = Meshes::Grid< 1, Real, Device, Index >; using RealType = typename GridType::RealType; using DeviceType = typename GridType::DeviceType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; constexpr static int getEntityDimension() { return 3; }; }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL