Skip to content

GitLab

Commit 6537a90c authored by Houman Mirzaalian Dastjerdi's avatar Houman Mirzaalian Dastjerdi
Browse files

Merge branches 'CUDA_ComplexGeometry' and 'List_Stability_checker'

parents b1c891c2 e1bfee69 67700050
Expand all Hide whitespace changes
Inline Side-by-side
Showing with 2429 additions and 257 deletions
.gitlab-ci.yml
View file @ 6537a90c
This diff is collapsed.
CMakeLists.txt
View file @ 6537a90c
......@@ -1027,31 +1027,22 @@ endif()
option ( WALBERLA_THREAD_SAFE_LOGGING "Enables/Disables thread-safe logging" ON )
if ( WALBERLA_BUILD_WITH_OPENMP )
if ( WALBERLA_CXX_COMPILER_IS_INTEL )
if ( WALBERLA_CXX_COMPILER_IS_INTEL AND "${CMAKE_CXX_COMPILER_VERSION}" VERSION_LESS "16.0.3" )
add_flag ( CMAKE_C_FLAGS "-openmp" )
add_flag ( CMAKE_CXX_FLAGS "-openmp" )
elseif ( CMAKE_COMPILER_IS_GNUCXX )
add_flag ( CMAKE_C_FLAGS "-fopenmp" )
add_flag ( CMAKE_CXX_FLAGS "-fopenmp" )
elseif ( WALBERLA_CXX_COMPILER_IS_CLANG )
add_flag ( CMAKE_C_FLAGS "-fopenmp" )
add_flag ( CMAKE_CXX_FLAGS "-fopenmp" )
elseif ( WALBERLA_CXX_COMPILER_IS_MSVC )
add_flag ( CMAKE_C_FLAGS "/openmp" )
add_flag ( CMAKE_CXX_FLAGS "/openmp" )
elseif ( WALBERLA_CXX_COMPILER_IS_IBM )
add_flag ( CMAKE_C_FLAGS "-qsmp=omp" )
add_flag ( CMAKE_CXX_FLAGS "-qsmp=omp" )
# There has been an internal compiler error with the IBM compiler, so WALBERLA_THREAD_SAFE_LOGGING is disabled by default for this compiler
set ( WALBERLA_THREAD_SAFE_LOGGING OFF CACHE BOOL "Enables/Disables thread-safe logging" FORCE )
elseif ( WALBERLA_CXX_COMPILER_IS_NEC )
add_flag ( CMAKE_C_FLAGS "-Popenmp" )
add_flag ( CMAKE_CXX_FLAGS "-Popenmp" )
add_flag ( CMAKE_C_FLAGS "-Popenmp" )
add_flag ( CMAKE_CXX_FLAGS "-Popenmp" )
else()
find_package( OpenMP )
add_flag ( CMAKE_C_FLAGS "${OpenMP_C_FLAGS}" )
add_flag ( CMAKE_CXX_FLAGS "${OpenMP_CXX_FLAGS}" )
list ( APPEND SERVICE_LIBS ${OpenMP_CXX_LIBRARIES} )
endif()
else()
if ( WALBERLA_CXX_COMPILER_IS_CRAY )
add_flag ( CMAKE_C_FLAGS "-h noomp" )
add_flag ( CMAKE_CXX_FLAGS "-h noomp" )
add_flag ( CMAKE_C_FLAGS "-h noomp" )
add_flag ( CMAKE_CXX_FLAGS "-h noomp" )
endif()
endif()
############################################################################################################################
......
apps/benchmarks/CMakeLists.txt
View file @ 6537a90c
add_subdirectory( ComplexGeometry )
add_subdirectory( ComplexGeometryList )
add_subdirectory( CUDA_ComplexGeometryList )
add_subdirectory( MeshDistance )
add_subdirectory( CouetteFlow )
add_subdirectory( ForcesOnSphereNearPlaneInShearFlow )
add_subdirectory( NonUniformGrid )
add_subdirectory( MotionSingleHeavySphere )
add_subdirectory( PeriodicGranularGas )
add_subdirectory( PoiseuilleChannel )
add_subdirectory( SchaeferTurek )
add_subdirectory( UniformGrid )
\ No newline at end of file
add_subdirectory( UniformGrid )
apps/benchmarks/CUDA_ComplexGeometryList/CMakeLists.txt 0 → 100644
View file @ 6537a90c
execute_process( COMMAND ${CMAKE_COMMAND} -E create_symlink
${CMAKE_CURRENT_SOURCE_DIR}/input.dat
${CMAKE_CURRENT_BINARY_DIR}/input.dat )
waLBerla_add_executable ( NAME CUDA_ComplexGeometryList
FILES CUDA_ComplexGeometryList.cpp CUDA_ComplexGeometryList.h
DEPENDS blockforest boundary core domain_decomposition field geometry mesh lbm postprocessing timeloop vtk simd cuda gui)
apps/benchmarks/CUDA_ComplexGeometryList/CUDA_ComplexGeometryList.cpp 0 → 100644
View file @ 6537a90c
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \file ComplexGeometryList.cpp
//! \author Christian Godenschwager <christian.godenschwager@fau.de>
//
//======================================================================================================================
#include "blockforest/communication/UniformBufferedScheme.h"
#include "core/Abort.h"
#include "core/Environment.h"
#include "core/config/Config.h"
#include "core/mpi/MPIManager.h"
#include "core/timing/RemainingTimeLogger.h"
#include "lbm/lattice_model/CollisionModel.h"
#include "lbm/lattice_model/D3Q19.h"
#include "lbm/list/CellCounters.h"
#include "lbm/list/List.h"
#include "lbm/list/ListCommunication.h"
#include "lbm/list/ListKernel.h"
#include "lbm/list/ListPressureBoundary.h"
#include "lbm/list/ListVTK.h"
#include "lbm/PerformanceEvaluation.h"
#include "lbm/BlockForestEvaluation.h"
#include "mesh/TriangleMeshes.h"
#include "mesh/MeshOperations.h"
#include "mesh/DistanceComputations.h"
#include "mesh/MeshIO.h"
#include "mesh/MatrixVectorOperations.h"
#include "mesh/blockforest/BlockForestInitialization.h"
#include "mesh/blockforest/BlockWorkloadMemory.h"
#include "mesh/blockforest/BlockExclusion.h"
#include "mesh/blockforest/RefinementSelection.h"
#include "mesh/distance_octree/DistanceOctree.h"
#include "mesh/boundary/BoundarySetup.h"
#include "mesh/boundary/BoundaryInfo.h"
#include "mesh/boundary/BoundaryLocation.h"
#include "mesh/boundary/BoundaryUIDFaceDataSource.h"
#include "mesh/boundary/ColorToBoundaryMapper.h"
#include "mesh/vtk/VTKMeshWriter.h"
#include "mesh/vtk/CommonDataSources.h"
#include "timeloop/SweepTimeloop.h"
#include "vtk/Initialization.h"
#include "vtk/VTKOutput.h"
#include "CUDA_ComplexGeometryList.h"
namespace complex_geometry{
using namespace walberla;
using namespace cuda;
template< typename MeshType >
void vertexToFaceColor( MeshType & mesh, const typename MeshType::Color & defaultColor )
{
WALBERLA_CHECK( mesh.has_vertex_colors() );
mesh.request_face_colors();
for( auto faceIt = mesh.faces_begin(); faceIt != mesh.faces_end(); ++faceIt )
{
typename MeshType::Color vertexColor;
bool useVertexColor = true;
auto vertexIt = mesh.fv_iter( *faceIt );
WALBERLA_ASSERT( vertexIt.is_valid() );
vertexColor = mesh.color( *vertexIt );
++vertexIt;
while( vertexIt.is_valid() && useVertexColor )
{
if( vertexColor != mesh.color( *vertexIt ) )
useVertexColor = false;
++vertexIt;
}
mesh.set_color( *faceIt, useVertexColor ? vertexColor : defaultColor );
}
}
template< typename MeshDistanceType, typename MeshType >
struct BoundaryLocationFunction
{
BoundaryLocationFunction( const shared_ptr< MeshDistanceType > & meshDistanceObject, const shared_ptr< mesh::BoundaryLocation< MeshType > > & boundaryLocation )
: meshDistanceObject_( meshDistanceObject ), boundaryLocation_( boundaryLocation ) { }
inline const mesh::BoundaryInfo & operator()( const Vector3< real_t > & p ) const
{
typename MeshType::FaceHandle fh;
meshDistanceObject_->sqSignedDistance( mesh::toOpenMesh( p ), fh );
return (*boundaryLocation_)[ fh ];
}
shared_ptr< MeshDistanceType > meshDistanceObject_;
shared_ptr< mesh::BoundaryLocation< MeshType > > boundaryLocation_;
};
template< typename MeshDistanceType, typename MeshType >
inline BoundaryLocationFunction< MeshDistanceType, MeshType > makeBoundaryLocationFunction( const shared_ptr< MeshDistanceType > & meshDistanceObject, const shared_ptr< mesh::BoundaryLocation< MeshType > > & boundaryLocation )
{
return BoundaryLocationFunction< MeshDistanceType, MeshType >( meshDistanceObject, boundaryLocation );
}
template< typename LBMList >
class MyLBMListVTKOutput {
public:
MyLBMListVTKOutput( const ConstBlockDataID & pdfList ) : pdfList_( pdfList ) {}
void operator()( std::vector< shared_ptr<vtk::BlockCellDataWriterInterface> > & writers,
std::map< std::string, vtk::VTKOutput::CellFilter > & filters,
std::map< std::string, vtk::VTKOutput::BeforeFunction > & /*beforeFunctions*/ )
{
writers.push_back( make_shared< lbm::ListVelocityVTKWriter< LBMList, float > >( pdfList_, "VelocityFromPDF" ) );
writers.push_back( make_shared< lbm::ListDensityVTKWriter < LBMList, float > >( pdfList_, "DensityFromPDF" ) );
writers.push_back( make_shared< lbm::ListPDFVTKWriter < LBMList, float > > ( pdfList_, "PDFs" ) );
filters["FluidFilter"] = lbm::ListFluidFilter<LBMList>( pdfList_ );
}
private:
const ConstBlockDataID pdfList_;
}; // class MyLBMListVTKOutput
template< typename MeshDistanceType >
struct MeshDistanceFunction
{
MeshDistanceFunction( const shared_ptr< MeshDistanceType > & meshDistanceObject ) : meshDistanceObject_( meshDistanceObject ) { }
inline real_t operator()( const Vector3< real_t > & p ) const { return real_c( meshDistanceObject_->sqSignedDistance( mesh::toOpenMesh( p ) ) ); }
shared_ptr< MeshDistanceType > meshDistanceObject_;
};
template< typename MeshDistanceType >
inline MeshDistanceFunction< MeshDistanceType > makeMeshDistanceFunction( const shared_ptr< MeshDistanceType > & meshDistanceObject )
{
return MeshDistanceFunction< MeshDistanceType >( meshDistanceObject );
}
int main( int argc, char **argv )
{
Environment env( argc, argv );
mpi::MPIManager::instance()->useWorldComm();
WALBERLA_CHECK_NOT_NULLPTR( env.config(), "You have to specify a config file!" );
Config::BlockHandle configBlock = env.config()->getBlock( "ComplexGeometry" );
if( !configBlock )
WALBERLA_ABORT( "You have to specify a \"ComplexGeometry\" block in the configuration file!" );
const real_t omega = configBlock.getParameter< real_t >( "omega", real_t(1.4) );
const real_t dx = configBlock.getParameter< real_t >( "dx" );
const Vector3<uint_t> blockSize = configBlock.getParameter< Vector3<uint_t> >( "blockSize" );
const uint_t timeSteps = configBlock.getParameter< uint_t >( "timeSteps" );
const std::string meshFile = configBlock.getParameter< std::string >( "meshFile" );
const real_t deltaPressure = configBlock.getParameter< real_t >( "deltaPressure" );
const uint_t perfLoggerInterval = configBlock.getParameter< uint_t >( "perfLoggerInterval", uint_t( 100 ) );
// Setup LatticeModel
typedef lbm::D3Q19< lbm::collision_model::TRT, false > MyLatticeModel;
MyLatticeModel latticeModel( lbm::collision_model::TRT::constructWithMagicNumber( omega ) );
// Load mesh and build DistanceOctree
auto mesh = make_shared< mesh::TriangleMesh >();
mesh->request_vertex_colors();
WALBERLA_LOG_DEVEL_ON_ROOT( "Loading mesh" );
mesh::readAndBroadcast( meshFile, *mesh);
vertexToFaceColor( *mesh, mesh::TriangleMesh::Color(255,255,255) );
WALBERLA_LOG_DEVEL_ON_ROOT( "Adding distance info to mesh" );
auto triDist = make_shared< mesh::TriangleDistance<mesh::TriangleMesh> >( mesh );
WALBERLA_LOG_DEVEL_ON_ROOT( "Building distance octree" );
auto distanceOctree = make_shared< mesh::DistanceOctree<mesh::TriangleMesh> >( triDist );
WALBERLA_LOG_DEVEL_ON_ROOT( "done. Octree has height " << distanceOctree->height() );
// Setup Blockforest
mesh::ComplexGeometryStructuredBlockforestCreator bfc( computeAABB( *mesh ), Vector3<real_t>( dx ), mesh::makeExcludeMeshInterior( distanceOctree, dx ) );
auto meshWorkloadMemory = mesh::makeMeshWorkloadMemory( distanceOctree, dx );
meshWorkloadMemory.setInsideCellWorkload(1);
meshWorkloadMemory.setOutsideCellWorkload(0);
bfc.setWorkloadMemorySUIDAssignmentFunction( meshWorkloadMemory );
auto blocks = bfc.createStructuredBlockForest( blockSize );
// Setup block data
typedef lbm::List< MyLatticeModel, lbm::LayoutSoA< MyLatticeModel::Stencil > > MyList;
BlockDataID pdfListId = lbm::addListToStorage< MyList >( blocks, "LBM list (FIdx)", latticeModel );
BlockDataID inflowPressureBoundaryHandling = lbm::addListPressureBoundaryToStorage< MyList >( pdfListId, blocks, "inflowBoundaryHandling" );
BlockDataID outflowPressureBoundaryHandling = lbm::addListPressureBoundaryToStorage< MyList >( pdfListId, blocks, "outflowBoundaryHandling" );
// CUDA
BlockDataID gpu_PdfListSrcID = addGPUListToStorage< MyList, ListOnGpu <uint32_t> >( blocks, pdfListId, "GPU Field Src" );
copyToGpu<MyList, ListOnGpu <uint32_t> >(*blocks, pdfListId, gpu_PdfListSrcID);
// Initialize block data
mesh::BoundarySetup bs( blocks, makeMeshDistanceFunction( distanceOctree ), uint_t(1) );
BoundaryUID noslipUID("noslip");
BoundaryUID inflowUID("inflowBoundaryHandling");
BoundaryUID outflowUID("outflowBoundaryHandling");
// attach boundary infos to mesh
mesh::ColorToBoundaryMapper<mesh::TriangleMesh> colorToBoundryMapper(( mesh::BoundaryInfo( noslipUID ) ));
colorToBoundryMapper.set( mesh::TriangleMesh::Color(255,0,0), mesh::BoundaryInfo( inflowUID ) );
colorToBoundryMapper.set( mesh::TriangleMesh::Color(0,0,255), mesh::BoundaryInfo( outflowUID ) );
colorToBoundryMapper.set( mesh::TriangleMesh::Color(255,255,255), mesh::BoundaryInfo( noslipUID ) );
auto boundaryLocations = colorToBoundryMapper.addBoundaryInfoToMesh( *mesh );
for( auto & block : *blocks )
{
MyList * lbmList = block.getData< MyList >( pdfListId );
lbm::ListPressureBoundary< MyList > * inflowPressureBH = block.getData< lbm::ListPressureBoundary< MyList > >( inflowPressureBoundaryHandling );
lbm::ListPressureBoundary< MyList > * outflowPressureBH = block.getData< lbm::ListPressureBoundary< MyList > >( outflowPressureBoundaryHandling );
WALBERLA_CHECK_NOT_NULLPTR( lbmList );
WALBERLA_CHECK_NOT_NULLPTR( inflowPressureBH );
WALBERLA_CHECK_NOT_NULLPTR( outflowPressureBH );
std::vector<Cell> fluidCells = bs.getDomainCells( block, mesh::BoundarySetup::INSIDE );
std::map< boundary::BoundaryUID, std::vector<Cell> > boundaryCells = bs.getBoundaryCells( block, mesh::BoundarySetup::OUTSIDE, makeBoundaryLocationFunction( distanceOctree, boundaryLocations ) );
lbmList->init( fluidCells );
inflowPressureBH->init( lbmList, boundaryCells[ inflowUID ], real_t( 1 ) + deltaPressure / real_t( 2 ) );
outflowPressureBH->init( lbmList, boundaryCells[ outflowUID ], real_t( 1 ) - deltaPressure / real_t( 2 ) );
}
// Blockforest output
lbm::BlockForestEvaluationBase< lbm::ListCellCounter, lbm::ListFluidCellCounter< MyList > > blockForestEval( blocks, lbm::ListCellCounter( blocks ),
lbm::ListFluidCellCounter< MyList >(blocks, pdfListId ) );
blockForestEval.logResultOnRoot();
// Setup Timeloop
SweepTimeloop timeloop( blocks->getBlockStorage(), timeSteps );
// Communication & Boundary Handling
blockforest::communication::UniformBufferedScheme< typename MyLatticeModel::Stencil > listScheme( blocks );
listScheme.addPackInfo( make_shared< lbm::ListPackInfo< MyList > >( pdfListId, blocks ) );
timeloop.add() << BeforeFunction( listScheme.getStartCommunicateFunctor(), "list communication start" )
<< Sweep( lbm::ListPressureBoundaryHandling< MyList >( inflowPressureBoundaryHandling ), "PressureBoundaryInflow" );
timeloop.add() << Sweep( lbm::ListPressureBoundaryHandling< MyList >( outflowPressureBoundaryHandling ), "PressureBoundaryOutflow" )
<< AfterFunction( listScheme.getWaitFunctor(), "list communication wait" );
// LBM Sweep
timeloop.add() << Sweep( lbm::ListDefaultTRTSweep< MyList >( pdfListId ), "ListSIMD2SplitTRTSweep" );
//timeloop.add() << Sweep( lbm::ListSIMD2SplitTRTSweep< MyList, 128 >( pdfListId ), "ListSIMD2SplitTRTSweep" );
// VTK Output
MyLBMListVTKOutput< MyList > myVTKOutput( pdfListId );
std::map< std::string, vtk::SelectableOutputFunction > vtkOutputFunctions;
vtk::initializeVTKOutput( vtkOutputFunctions, myVTKOutput, blocks, env.config() );
for( auto output = vtkOutputFunctions.begin(); output != vtkOutputFunctions.end(); ++output )
timeloop.addFuncBeforeTimeStep( output->second.outputFunction, std::string( "VTK: " ) + output->first,
output->second.requiredGlobalStates, output->second.incompatibleGlobalStates );
// Utilities
lbm::PerformanceEvaluationBase< lbm::ListCellCounter, lbm::ListFluidCellCounter< MyList > > performanceEval( blocks, lbm::ListCellCounter( blocks ),
lbm::ListFluidCellCounter< MyList >(blocks, pdfListId ) );
WcTimer performanceTimer;
timeloop.addFuncAfterTimeStep( [&]() {
if( timeloop.getCurrentTimeStep() == 0 )
{
performanceTimer.start();
}
else if( timeloop.getCurrentTimeStep() % perfLoggerInterval == 0 )
{
performanceTimer.end();
performanceEval.logResultOnRoot( perfLoggerInterval, performanceTimer.last() );
performanceTimer.start();
}
}, "PerformanceEvaluation");
timeloop.addFuncAfterTimeStep( timing::RemainingTimeLogger( timeloop.getNrOfTimeSteps() ) );
// Run timeloop
WcTimingPool timingPool;
timeloop.run( timingPool );
timingPool.logResultOnRoot( walberla::WcTimingPool::REDUCE_TOTAL, true );
return EXIT_SUCCESS;
}
} // namespace complex_geometry
int main( int argc, char ** argv )
{
return complex_geometry::main( argc, argv );
}
apps/benchmarks/CUDA_ComplexGeometryList/CUDA_ComplexGeometryList.cu 0 → 100644
View file @ 6537a90c
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \file AddGPUListToStorage.h
//! \ingroup cuda
//! \author Houman Mirzaalian D. <houman.mirzaalian@fau.de>
//
//======================================================================================================================
#include "CUDA_ComplexGeometryList.h"
apps/benchmarks/CUDA_ComplexGeometryList/CUDA_ComplexGeometryList.h 0 → 100644
View file @ 6537a90c
//======================================================================================================================
//
// This file is part of waLBerla. waLBerla is free software: you can
// redistribute it and/or modify it under the terms of the GNU General Public
// License as published by the Free Software Foundation, either version 3 of
// the License, or (at your option) any later version.
//
// waLBerla is distributed in the hope that it will be useful, but WITHOUT
// ANY WARRANTY; without even the implied warranty of MERCHANTABILITY or
// FITNESS FOR A PARTICULAR PURPOSE. See the GNU General Public License
// for more details.
//
// You should have received a copy of the GNU General Public License along
// with waLBerla (see COPYING.txt). If not, see <http://www.gnu.org/licenses/>.
//
//! \file AddGPUListToStorage.h
//! \ingroup cuda
//! \author Houman Mirzaalian D. <houman.mirzaalian@fau.de>
//
//======================================================================================================================
#pragma once
#include "lbm/list/List.h"
#include "lbm/all.h"
#include <lbm/lattice_model/LatticeModelBase.h>
#include "lbm/list/List.h"
#include "lbm/lattice_model/D3Q19.h"
#include "domain_decomposition/BlockDataID.h"
#include "cuda/HostFieldAllocator.h"
#include "cuda/HostFieldAllocator.h"
#include "cuda/FieldCopy.h"
#include "cuda/GPUField.h"
#include "cuda/Kernel.h"
#include "cuda/AddGPUFieldToStorage.h"
#include "cuda/communication/GPUPackInfo.h"
#include "cuda/FieldIndexing.h"
#include "CUDA_ComplexGeometryList.h"
#include <type_traits>
namespace walberla {
namespace cuda {
using lbm::List;
using lbm::LayoutSoA;
typedef lbm::D3Q19<lbm::collision_model::TRT, false> LatticeModel_T;
template<typename Index_T = walberla::uint32_t>
class ListOnGpu {
public:
typedef Index_T index_t;
ListOnGpu() : pdfs_(nullptr), tmpPdfs_(nullptr), pullIdxs_(nullptr), size_(0), numFluidCells_(0) {}
~ListOnGpu() {};
void resize(const size_t newSize, const uint_t newNumFluidCells) {
cudaFree(pdfs_);
cudaFree(tmpPdfs_);
cudaFree(pullIdxs_);
cudaMalloc(&pdfs_, sizeof(real_t) * newSize);
cudaMalloc(&tmpPdfs_, sizeof(real_t) * newSize);
cudaMalloc(&pullIdxs_, sizeof(Index_T) * newSize);
size_ = newSize;
numFluidCells_ = newNumFluidCells;
//free old ptr and reallocate
}
real_t *getPdfData() { return pdfs_; }
const real_t *getPdfData() const { return pdfs_; }
real_t *getTmpPdfData() { return tmpPdfs_; }
const real_t *getTmpPdfData() const { return tmpPdfs_; }
Index_T *getPullIndicesData() { return pullIdxs_; }
const Index_T *getPullIndicesData() const { return pullIdxs_; }
size_t getSize() const { return size_; }
size_t getNumFluidCells() const { return numFluidCells_; }
private:
real_t *pdfs_;
real_t *tmpPdfs_;
Index_T *pullIdxs_;
size_t size_;
uint_t numFluidCells_;
};
template< typename Index_T >
bool operator==( const ListOnGpu<Index_T> & lhs, const ListOnGpu<Index_T> & rhs )
{
return false;
}
template< typename ListOnCPU, typename ListOnGPU >
void copyToGpu( const ListOnCPU & src, ListOnGPU & dst )
{
static_assert( std::is_same< typename ListOnCPU::index_t, typename ListOnGPU::index_t >::value, "" );
typedef typename ListOnCPU::index_t index_t;
auto pdfs_cpu = src.getPdfData();
auto pullIdxs_cpu = src.getPullIndicesData();
auto size_cpu = src.getSize();
auto size_FluidCells = src.getNumFluidCells();
auto pdfs_gpu = dst.getPdfData();
auto pullIdxs_gpu = dst.getPullIndicesData();
auto size_gpu = dst.getSize();
if (size_cpu != size_gpu ){
dst.resize(size_cpu, size_FluidCells);
} // check if size matches, if not reallocate
cudaMemcpy(pdfs_gpu, pdfs_cpu, sizeof(real_t) * size_cpu, cudaMemcpyHostToDevice);
cudaMemcpy(pullIdxs_gpu, pullIdxs_cpu, sizeof(index_t) * size_cpu, cudaMemcpyHostToDevice);
// memcopy
}
template< typename ListOnGPU, typename ListOnCPU>
void copyToCpu( const ListOnGPU & src, ListOnCPU & dst )
{
static_assert( std::is_same< typename ListOnCPU::index_t, typename ListOnGPU::index_t >::value, "" );