Loading src/Benchmarks/Traversers/GridTraverserBenchmarkHelper.h +4 −132 Original line number Diff line number Diff line Loading @@ -20,143 +20,15 @@ namespace TNL { namespace Benchmarks { namespace Traversers { #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; 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(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( 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 ); } } #endif template< typename Grid, typename Device = typename Grid::DeviceType > class GridTraverserBenchmarkHelper{}; template< typename Grid > class GridTraverserBenchmarkHelper< Grid, Devices::Host > { public: using GridType = Grid; using GridPointer = Pointers::SharedPointer< Grid >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename Grid::CoordinatesType; using MeshFunction = Functions::MeshFunction< Grid >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename Grid::template EntityType< 1, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< Grid, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void noBCTraverserTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { /*Meshes::GridTraverser< Grid >::template processEntities< CellType, AddOneEntitiesProcessorType, UserDataType, false >( grid, CoordinatesType( 0 ), grid->getDimensions() - CoordinatesType( 1 ), userData ); */ const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); //MeshFunction* _u = &u.template modifyData< Device >(); SimpleCellType entity( *grid ); for( IndexType x = begin.x(); x <= end.x(); x ++ ) { entity.getCoordinates().x() = x; entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); } } }; template< typename Grid > class GridTraverserBenchmarkHelper< Grid, Devices::Cuda > { public: using GridType = Grid; using GridPointer = Pointers::SharedPointer< Grid >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename Grid::CoordinatesType; using MeshFunction = Functions::MeshFunction< Grid >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename Grid::template EntityType< 1, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< Grid, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void noBCTraverserTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 256 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size ); dim3 gridIdx; for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser1D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); class GridTraverserBenchmarkHelper{}; } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL #include "GridTraverserBenchmarkHelper_1D.h" #include "GridTraverserBenchmarkHelper_2D.h" #include "GridTraverserBenchmarkHelper_3D.h" src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_1D.h 0 → 100644 +154 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_1D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #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; 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(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( 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 ); } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 1, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 1; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 1, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 1; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 256 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size ); dim3 gridIdx; for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser1D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_2D.h 0 → 100644 +152 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_2D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #ifdef HAVE_CUDA template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor > __global__ void _GridTraverser2D( const Meshes::Grid< 2, Real, Devices::Cuda, Index >* grid, UserData userData, const typename GridEntity::CoordinatesType begin, const typename GridEntity::CoordinatesType end, const dim3 gridIdx ) { typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 2, Real, Devices::Cuda, Index > GridType; GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx.x * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; entity.getCoordinates().y() = begin.y() + ( gridIdx.y * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y; if( entity.getCoordinates() <= end ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 2, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 2; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().y() = begin.y(); entity.getCoordinates().y() <= end.y(); entity.getCoordinates().y()++ ) for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 2, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 2; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 16, 16 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size, size ); dim3 gridIdx; for( gridIdx.y = 0; gridIdx.y < gridsCount.y; gridIdx.y++ ) for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser2D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_3D.h 0 → 100644 +156 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_3D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #ifdef HAVE_CUDA template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor > __global__ void _GridTraverser3D( const Meshes::Grid< 3, Real, Devices::Cuda, Index >* grid, UserData userData, const typename GridEntity::CoordinatesType begin, const typename GridEntity::CoordinatesType end, const dim3 gridIdx ) { typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 3, Real, Devices::Cuda, Index > GridType; GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx.x * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; entity.getCoordinates().y() = begin.y() + ( gridIdx.y * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y; entity.getCoordinates().z() = begin.z() + ( gridIdx.z * Devices::Cuda::getMaxGridSize() + blockIdx.z ) * blockDim.z + threadIdx.z; if( entity.getCoordinates() <= end ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 3, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 3; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().z() = begin.z(); entity.getCoordinates().z() <= end.z(); entity.getCoordinates().z()++ ) for( entity.getCoordinates().y() = begin.y(); entity.getCoordinates().y() <= end.y(); entity.getCoordinates().y()++ ) for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 3, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 3; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 32, 4, 2 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size, size, size ); dim3 gridIdx; for( gridIdx.z = 0; gridIdx.z < gridsCount.z; gridIdx.z++ ) for( gridIdx.y = 0; gridIdx.y < gridsCount.y; gridIdx.y++ ) for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser3D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL src/Benchmarks/Traversers/GridTraversersBenchmark_1D.h +25 −19 File changed.Preview size limit exceeded, changes collapsed. Show changes Loading
src/Benchmarks/Traversers/GridTraverserBenchmarkHelper.h +4 −132 Original line number Diff line number Diff line Loading @@ -20,143 +20,15 @@ namespace TNL { namespace Benchmarks { namespace Traversers { #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; 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(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( 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 ); } } #endif template< typename Grid, typename Device = typename Grid::DeviceType > class GridTraverserBenchmarkHelper{}; template< typename Grid > class GridTraverserBenchmarkHelper< Grid, Devices::Host > { public: using GridType = Grid; using GridPointer = Pointers::SharedPointer< Grid >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename Grid::CoordinatesType; using MeshFunction = Functions::MeshFunction< Grid >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename Grid::template EntityType< 1, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< Grid, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void noBCTraverserTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { /*Meshes::GridTraverser< Grid >::template processEntities< CellType, AddOneEntitiesProcessorType, UserDataType, false >( grid, CoordinatesType( 0 ), grid->getDimensions() - CoordinatesType( 1 ), userData ); */ const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); //MeshFunction* _u = &u.template modifyData< Device >(); SimpleCellType entity( *grid ); for( IndexType x = begin.x(); x <= end.x(); x ++ ) { entity.getCoordinates().x() = x; entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; //AddOneEntitiesProcessorType::processEntity( entity.getMesh(), userData, entity ); } } }; template< typename Grid > class GridTraverserBenchmarkHelper< Grid, Devices::Cuda > { public: using GridType = Grid; using GridPointer = Pointers::SharedPointer< Grid >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename Grid::CoordinatesType; using MeshFunction = Functions::MeshFunction< Grid >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename Grid::template EntityType< 1, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< Grid, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void noBCTraverserTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 256 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size ); dim3 gridIdx; for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser1D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); class GridTraverserBenchmarkHelper{}; } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL #include "GridTraverserBenchmarkHelper_1D.h" #include "GridTraverserBenchmarkHelper_2D.h" #include "GridTraverserBenchmarkHelper_3D.h"
src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_1D.h 0 → 100644 +154 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_1D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #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; 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(); //( userData.u->getData() )[ entity.getIndex( coordinates ) ] += ( RealType ) 1.0; //( userData.u->getData() )[ coordinates.x() ] += ( 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 ); } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 1, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 1; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 1, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 1; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 256 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size ); dim3 gridIdx; for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser1D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL
src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_2D.h 0 → 100644 +152 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_2D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #ifdef HAVE_CUDA template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor > __global__ void _GridTraverser2D( const Meshes::Grid< 2, Real, Devices::Cuda, Index >* grid, UserData userData, const typename GridEntity::CoordinatesType begin, const typename GridEntity::CoordinatesType end, const dim3 gridIdx ) { typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 2, Real, Devices::Cuda, Index > GridType; GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx.x * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; entity.getCoordinates().y() = begin.y() + ( gridIdx.y * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y; if( entity.getCoordinates() <= end ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 2, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 2; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().y() = begin.y(); entity.getCoordinates().y() <= end.y(); entity.getCoordinates().y()++ ) for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); //userData.u->getData()[ entity.getIndex() ] += ( RealType ) 1.0; ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 2, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 2; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 16, 16 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size, size ); dim3 gridIdx; for( gridIdx.y = 0; gridIdx.y < gridsCount.y; gridIdx.y++ ) for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser2D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL
src/Benchmarks/Traversers/GridTraverserBenchmarkHelper_3D.h 0 → 100644 +156 −0 Original line number Diff line number Diff line /*************************************************************************** GridTraversersBenchmarkHelper_3D.h - description ------------------- begin : Jan 6, 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 "GridTraverserBenchmarkHelper.h" #include "AddOneEntitiesProcessor.h" #include "BenchmarkTraverserUserData.h" #include "SimpleCell.h" namespace TNL { namespace Benchmarks { namespace Traversers { #ifdef HAVE_CUDA template< typename Real, typename Index, typename GridEntity, typename UserData, typename EntitiesProcessor > __global__ void _GridTraverser3D( const Meshes::Grid< 3, Real, Devices::Cuda, Index >* grid, UserData userData, const typename GridEntity::CoordinatesType begin, const typename GridEntity::CoordinatesType end, const dim3 gridIdx ) { typedef Real RealType; typedef Index IndexType; typedef Meshes::Grid< 3, Real, Devices::Cuda, Index > GridType; GridEntity entity( *grid ); entity.getCoordinates().x() = begin.x() + ( gridIdx.x * Devices::Cuda::getMaxGridSize() + blockIdx.x ) * blockDim.x + threadIdx.x; entity.getCoordinates().y() = begin.y() + ( gridIdx.y * Devices::Cuda::getMaxGridSize() + blockIdx.y ) * blockDim.y + threadIdx.y; entity.getCoordinates().z() = begin.z() + ( gridIdx.z * Devices::Cuda::getMaxGridSize() + blockIdx.z ) * blockDim.z + threadIdx.z; if( entity.getCoordinates() <= end ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } #endif template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 3, Real, Devices::Host, Index > > { public: constexpr static int Dimension = 3; using GridType = Meshes::Grid< Dimension, Real, Devices::Host, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { const CoordinatesType begin( 0 ); const CoordinatesType end = CoordinatesType( size ) - CoordinatesType( 1 ); SimpleCellType entity( *grid ); for( entity.getCoordinates().z() = begin.z(); entity.getCoordinates().z() <= end.z(); entity.getCoordinates().z()++ ) for( entity.getCoordinates().y() = begin.y(); entity.getCoordinates().y() <= end.y(); entity.getCoordinates().y()++ ) for( entity.getCoordinates().x() = begin.x(); entity.getCoordinates().x() <= end.x(); entity.getCoordinates().x() ++ ) { entity.refresh(); ( *userData.u )( entity ) += ( RealType ) 1.0; } } }; template< typename Real, typename Index > class GridTraverserBenchmarkHelper< Meshes::Grid< 3, Real, Devices::Cuda, Index > > { public: constexpr static int Dimension = 3; using GridType = Meshes::Grid< Dimension, Real, Devices::Cuda, Index >; using GridPointer = Pointers::SharedPointer< GridType >; using RealType = typename GridType::RealType; using IndexType = typename GridType::IndexType; using CoordinatesType = typename GridType::CoordinatesType; using MeshFunction = Functions::MeshFunction< GridType >; using MeshFunctionPointer = Pointers::SharedPointer< MeshFunction >; using CellType = typename GridType::template EntityType< Dimension, Meshes::GridEntityNoStencilStorage >; using SimpleCellType = SimpleCell< GridType >; using Traverser = Meshes::Traverser< GridType, CellType >; using UserDataType = BenchmarkTraverserUserData< MeshFunction >; using AddOneEntitiesProcessorType = AddOneEntitiesProcessor< UserDataType >; static void simpleCellTest( const GridPointer& grid, UserDataType& userData, std::size_t size ) { #ifdef HAVE_CUDA dim3 blockSize( 32, 4, 2 ), blocksCount, gridsCount; Devices::Cuda::setupThreads( blockSize, blocksCount, gridsCount, size, size, size ); dim3 gridIdx; for( gridIdx.z = 0; gridIdx.z < gridsCount.z; gridIdx.z++ ) for( gridIdx.y = 0; gridIdx.y < gridsCount.y; gridIdx.y++ ) for( gridIdx.x = 0; gridIdx.x < gridsCount.x; gridIdx.x++ ) { dim3 gridSize; Devices::Cuda::setupGrid( blocksCount, gridsCount, gridIdx, gridSize ); _GridTraverser3D< RealType, IndexType, SimpleCellType, UserDataType, AddOneEntitiesProcessorType > <<< blocksCount, blockSize >>> ( &grid.template getData< Devices::Cuda >(), userData, CoordinatesType( 0 ), CoordinatesType( size ) - CoordinatesType( 1 ), gridIdx.x ); } #endif } }; } // namespace Traversers } // namespace Benchmarks } // namespace TNL
src/Benchmarks/Traversers/GridTraversersBenchmark_1D.h +25 −19 File changed.Preview size limit exceeded, changes collapsed. Show changes