OilfieldToolsNet/opm-simulators
4
0
Fork 0
mirror of https://github.com/OPM/opm-simulators.git synced 2025-02-25 18:55:30 -06:00
Code Issues Packages Projects Releases Wiki Activity

added: option to disable the BDA solvers

Browse Source
This commit is contained in:
Arne Morten Kvarving 2023-08-09 15:39:32 +02:00 committed by hnil
parent 07fb18422d
commit 896cb8484d
8 changed files with 257 additions and 223 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

49
CMakeLists.txt
View File

@ -31,6 +31,7 @@ option(USE_CHOW_PATEL_ILU_GPU "Run iterative ILU decomposition on GPU? Requires
option(USE_CHOW_PATEL_ILU_GPU_PARALLEL "Try to use more parallelism on the GPU during the iterative ILU decomposition? Requires USE_CHOW_PATEL_ILU_GPU" OFF)
option(BUILD_FLOW_ALU_GRID "Build flow blackoil with alu grid" OFF)
option(USE_DAMARIS_LIB "Use the Damaris library for asynchronous I/O?" OFF)
option(USE_BDA_BRIDGE "Enable the BDA bridge (GPU/AMGCL solvers)" ON)
# The following was copied from CMakeLists.txt in opm-common.
# TODO: factor out the common parts in opm-common and opm-simulator as a cmake module
@ -102,6 +103,10 @@ if(USE_MPI)
set(HDF5_PREFER_PARALLEL TRUE)
endif()
if(USE_BDA_BRIDGE)
set(COMPILE_BDA_BRIDGE 1)
endif()
# not the same location as most of the other projects? this hook overrides
macro (dir_hook)
endmacro (dir_hook)
@ -170,7 +175,6 @@ endif()
if(CUDA_FOUND)
set(HAVE_CUDA 1)
set(COMPILE_BDA_BRIDGE 1)
include_directories(${CUDA_INCLUDE_DIRS})
endif()
@ -182,7 +186,6 @@ if(OpenCL_FOUND)
find_file(CL2_HPP CL/cl2.hpp HINTS ${OpenCL_INCLUDE_DIRS})
if(CL2_HPP)
set(HAVE_OPENCL 1)
set(COMPILE_BDA_BRIDGE 1)
include_directories(${OpenCL_INCLUDE_DIRS})
find_file(OPENCL_HPP CL/opencl.hpp HINTS ${OpenCL_INCLUDE_DIRS})
if(OPENCL_HPP)
@ -216,7 +219,6 @@ endif()
find_package(amgcl)
if(amgcl_FOUND)
set(HAVE_AMGCL 1)
set(COMPILE_BDA_BRIDGE 1)
# Linking to target angcl::amgcl drags in OpenMP and -fopenmp as a compile
# flag. With that nvcc fails as it does not that flag.
# Hence we set AMGCL_INCLUDE_DIRS.
@ -228,7 +230,6 @@ if(OpenCL_FOUND)
find_package(VexCL)
if(VexCL_FOUND)
set(HAVE_VEXCL 1)
set(COMPILE_BDA_BRIDGE 1)
# generator expressions in vexcl do not seem to work and therefore
# we cannot use the imported target. Hence we exract the needed info
# from the targets
@ -287,14 +288,12 @@ macro (files_hook)
else()
if(rocsparse_FOUND AND rocblas_FOUND)
set(HAVE_ROCSPARSE 1)
set(COMPILE_BDA_BRIDGE 1)
else()
unset(HAVE_ROCSPARSE)
endif()
endif()
if(ROCALUTION_FOUND)
set(HAVE_ROCALUTION 1)
set(COMPILE_BDA_BRIDGE 1)
endif()
endif()
if(MPI_FOUND AND HDF5_FOUND AND NOT HDF5_IS_PARALLEL)
@ -632,7 +631,9 @@ add_custom_target(extra_test ${CMAKE_CTEST_COMMAND} -C ExtraTests)
if(CUDA_FOUND)
target_link_libraries( opmsimulators PUBLIC ${CUDA_cusparse_LIBRARY} )
target_link_libraries( opmsimulators PUBLIC ${CUDA_cublas_LIBRARY} )
set_tests_properties(cusparseSolver PROPERTIES LABELS gpu_cuda)
if(USE_BDA_BRIDGE)
set_tests_properties(cusparseSolver PROPERTIES LABELS gpu_cuda)
endif()
# CUISTL
set_tests_properties(cusparse_safe_call
@ -649,25 +650,27 @@ if(CUDA_FOUND)
PROPERTIES LABELS gpu_cuda)
endif()
if(OpenCL_FOUND)
target_link_libraries( opmsimulators PUBLIC ${OpenCL_LIBRARIES} )
set_tests_properties(openclSolver solvetransposed3x3 csrToCscOffsetMap
PROPERTIES LABELS gpu_opencl)
endif()
if(USE_BDA_BRIDGE)
if(OpenCL_FOUND)
target_link_libraries( opmsimulators PUBLIC ${OpenCL_LIBRARIES} )
set_tests_properties(openclSolver solvetransposed3x3 csrToCscOffsetMap
PROPERTIES LABELS gpu_opencl)
endif()
if(ROCALUTION_FOUND)
target_include_directories(opmsimulators PUBLIC ${rocalution_INCLUDE_DIR}/rocalution)
set_tests_properties(rocalutionSolver PROPERTIES LABELS gpu_rocm)
endif()
if(ROCALUTION_FOUND)
target_include_directories(opmsimulators PUBLIC ${rocalution_INCLUDE_DIR}/rocalution)
set_tests_properties(rocalutionSolver PROPERTIES LABELS gpu_rocm)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
target_link_libraries( opmsimulators PUBLIC roc::rocsparse )
target_link_libraries( opmsimulators PUBLIC roc::rocblas )
set_tests_properties(rocsparseSolver PROPERTIES LABELS gpu_rocm)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
target_link_libraries( opmsimulators PUBLIC roc::rocsparse )
target_link_libraries( opmsimulators PUBLIC roc::rocblas )
set_tests_properties(rocsparseSolver PROPERTIES LABELS gpu_rocm)
endif()
if(VexCL_FOUND)
target_link_libraries( opmsimulators PUBLIC OPM::VexCL::OpenCL )
if(VexCL_FOUND)
target_link_libraries( opmsimulators PUBLIC OPM::VexCL::OpenCL )
endif()
endif()
if(Damaris_FOUND)

93
CMakeLists_files.cmake
View File

@ -52,8 +52,6 @@ list (APPEND MAIN_SOURCE_FILES
opm/simulators/flow/SimulatorFullyImplicitBlackoilEbos.cpp
opm/simulators/flow/ValidationFunctions.cpp
opm/simulators/flow/partitionCells.cpp
opm/simulators/linalg/bda/WellContributions.cpp
opm/simulators/linalg/bda/MultisegmentWellContribution.cpp
opm/simulators/linalg/ExtractParallelGridInformationToISTL.cpp
opm/simulators/linalg/FlexibleSolver1.cpp
opm/simulators/linalg/FlexibleSolver2.cpp
@ -62,7 +60,6 @@ list (APPEND MAIN_SOURCE_FILES
opm/simulators/linalg/FlexibleSolver5.cpp
opm/simulators/linalg/FlexibleSolver6.cpp
opm/simulators/linalg/ISTLSolverEbos.cpp
opm/simulators/linalg/ISTLSolverEbosWithGPU.cpp
opm/simulators/linalg/MILU.cpp
opm/simulators/linalg/ParallelIstlInformation.cpp
opm/simulators/linalg/ParallelOverlappingILU0.cpp
@ -143,14 +140,17 @@ list (APPEND MAIN_SOURCE_FILES
opm/simulators/wells/WGState.cpp
)
if (DAMARIS_FOUND AND MPI_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/utils/DamarisOutputModule.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/utils/DamarisKeywords.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/utils/initDamarisXmlFile.cpp)
endif()
if(CUDA_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/cuWellContributions.cu)
if(USE_BDA_BRIDGE)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/cuWellContributions.cu)
endif()
# CUISTL SOURCE
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/cuistl/detail/CuBlasHandle.cpp)
@ -191,33 +191,36 @@ if(CUDA_FOUND)
list (APPEND PUBLIC_HEADER_FILES opm/simulators/linalg/cuistl/set_device.hpp)
endif()
if(OPENCL_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/BlockedMatrix.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/BILU0.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/Reorder.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/ChowPatelIlu.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/BISAI.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/CPR.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/opencl.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclKernels.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/OpenclMatrix.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/Preconditioner.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclWellContributions.cpp)
endif()
if(ROCALUTION_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/rocalutionSolverBackend.cpp)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/rocsparseSolverBackend.cpp)
endif()
if(COMPILE_BDA_BRIDGE)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/BdaBridge.cpp)
endif()
if(amgcl_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/amgclSolverBackend.cpp)
if(CUDA_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/amgclSolverBackend.cu)
if(USE_BDA_BRIDGE)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/BdaBridge.cpp
opm/simulators/linalg/bda/WellContributions.cpp
opm/simulators/linalg/bda/MultisegmentWellContribution.cpp
opm/simulators/linalg/ISTLSolverEbosWithGPU.cpp)
if(OPENCL_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/BlockedMatrix.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/BILU0.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/Reorder.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/ChowPatelIlu.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/BISAI.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/CPR.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/opencl.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclKernels.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/OpenclMatrix.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/Preconditioner.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/opencl/openclWellContributions.cpp)
endif()
if(ROCALUTION_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/rocalutionSolverBackend.cpp)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/rocsparseSolverBackend.cpp)
endif()
if(amgcl_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/amgclSolverBackend.cpp)
if(CUDA_FOUND)
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cuda/amgclSolverBackend.cu)
endif()
endif()
endif()
if(MPI_FOUND)
@ -271,7 +274,9 @@ if(MPI_FOUND)
endif()
if(CUDA_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_cusparseSolver.cpp)
if(USE_BDA_BRIDGE)
list(APPEND TEST_SOURCE_FILES tests/test_cusparseSolver.cpp)
endif()
list(APPEND TEST_SOURCE_FILES tests/cuistl/test_cusparse_safe_call.cpp)
list(APPEND TEST_SOURCE_FILES tests/cuistl/test_cublas_safe_call.cpp)
list(APPEND TEST_SOURCE_FILES tests/cuistl/test_cuda_safe_call.cpp)
@ -288,17 +293,21 @@ if(CUDA_FOUND)
endif()
if(OPENCL_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_openclSolver.cpp)
list(APPEND TEST_SOURCE_FILES tests/test_solvetransposed3x3.cpp)
if(USE_BDA_BRIDGE)
if(OPENCL_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_openclSolver.cpp)
list(APPEND TEST_SOURCE_FILES tests/test_solvetransposed3x3.cpp)
list(APPEND TEST_SOURCE_FILES tests/test_csrToCscOffsetMap.cpp)
endif()
if(ROCALUTION_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_rocalutionSolver.cpp)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_rocsparseSolver.cpp)
endif()
endif()
if(ROCALUTION_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_rocalutionSolver.cpp)
endif()
if(rocsparse_FOUND AND rocblas_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_rocsparseSolver.cpp)
endif()
if(HDF5_FOUND)
list(APPEND TEST_SOURCE_FILES tests/test_HDF5File.cpp)
list(APPEND TEST_SOURCE_FILES tests/test_HDF5Serializer.cpp)

5
opm/simulators/flow/BlackoilModelEbosNldd.hpp
View File

@ -34,7 +34,12 @@
#include <opm/simulators/flow/SubDomain.hpp>
#include <opm/simulators/linalg/extractMatrix.hpp>
#if COMPILE_BDA_BRIDGE
#include <opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp>
#else
#include <opm/simulators/linalg/ISTLSolverEbos.hpp>
#endif
#include <opm/simulators/timestepping/ConvergenceReport.hpp>
#include <opm/simulators/timestepping/SimulatorReport.hpp>

312
opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp
View File

@ -21,14 +21,15 @@
#ifndef OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED
#define OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED
#include <opm/simulators/linalg/ISTLSolverEbos.hpp>
#if COMPILE_BDA_BRIDGE
namespace Opm
{
namespace Opm {
template<class Matrix, class Vector, int block_size> class BdaBridge;
class WellContributions;
namespace detail
{
namespace detail {
template<class Matrix, class Vector>
struct BdaSolverInfo
{
@ -83,179 +84,180 @@ private:
};
}
/// This class solves the fully implicit black-oil system by
/// solving the reduced system (after eliminating well variables)
/// as a block-structured matrix (one block for all cell variables) for a fixed
/// number of cell variables np .
template <class TypeTag>
class ISTLSolverEbosWithGPU : public ISTLSolverEbos<TypeTag>
{
protected:
using ParentType = ISTLSolverEbos<TypeTag>;
using GridView = GetPropType<TypeTag, Properties::GridView>;
using Scalar = GetPropType<TypeTag, Properties::Scalar>;
using SparseMatrixAdapter = GetPropType<TypeTag, Properties::SparseMatrixAdapter>;
using Vector = GetPropType<TypeTag, Properties::GlobalEqVector>;
using Indices = GetPropType<TypeTag, Properties::Indices>;
using WellModel = GetPropType<TypeTag, Properties::EclWellModel>;
using Simulator = GetPropType<TypeTag, Properties::Simulator>;
using Matrix = typename SparseMatrixAdapter::IstlMatrix;
using ThreadManager = GetPropType<TypeTag, Properties::ThreadManager>;
using ElementContext = GetPropType<TypeTag, Properties::ElementContext>;
using AbstractSolverType = Dune::InverseOperator<Vector, Vector>;
using AbstractOperatorType = Dune::AssembledLinearOperator<Matrix, Vector, Vector>;
using AbstractPreconditionerType = Dune::PreconditionerWithUpdate<Vector, Vector>;
using WellModelOperator = WellModelAsLinearOperator<WellModel, Vector, Vector>;
using ElementMapper = GetPropType<TypeTag, Properties::ElementMapper>;
constexpr static std::size_t pressureIndex = GetPropType<TypeTag, Properties::Indices>::pressureSwitchIdx;
/// This class solves the fully implicit black-oil system by
/// solving the reduced system (after eliminating well variables)
/// as a block-structured matrix (one block for all cell variables) for a fixed
/// number of cell variables np .
template <class TypeTag>
class ISTLSolverEbosWithGPU : public ISTLSolverEbos<TypeTag>
{
protected:
using ParentType = ISTLSolverEbos<TypeTag>;
using GridView = GetPropType<TypeTag, Properties::GridView>;
using Scalar = GetPropType<TypeTag, Properties::Scalar>;
using SparseMatrixAdapter = GetPropType<TypeTag, Properties::SparseMatrixAdapter>;
using Vector = GetPropType<TypeTag, Properties::GlobalEqVector>;
using Indices = GetPropType<TypeTag, Properties::Indices>;
using WellModel = GetPropType<TypeTag, Properties::EclWellModel>;
using Simulator = GetPropType<TypeTag, Properties::Simulator>;
using Matrix = typename SparseMatrixAdapter::IstlMatrix;
using ThreadManager = GetPropType<TypeTag, Properties::ThreadManager>;
using ElementContext = GetPropType<TypeTag, Properties::ElementContext>;
using AbstractSolverType = Dune::InverseOperator<Vector, Vector>;
using AbstractOperatorType = Dune::AssembledLinearOperator<Matrix, Vector, Vector>;
using AbstractPreconditionerType = Dune::PreconditionerWithUpdate<Vector, Vector>;
using WellModelOperator = WellModelAsLinearOperator<WellModel, Vector, Vector>;
using ElementMapper = GetPropType<TypeTag, Properties::ElementMapper>;
constexpr static std::size_t pressureIndex = GetPropType<TypeTag, Properties::Indices>::pressureSwitchIdx;
#if HAVE_MPI
using CommunicationType = Dune::OwnerOverlapCopyCommunication<int,int>;
using CommunicationType = Dune::OwnerOverlapCopyCommunication<int,int>;
#else
using CommunicationType = Dune::CollectiveCommunication<int>;
using CommunicationType = Dune::CollectiveCommunication<int>;
#endif
public:
using AssembledLinearOperatorType = Dune::AssembledLinearOperator< Matrix, Vector, Vector >;
public:
using AssembledLinearOperatorType = Dune::AssembledLinearOperator< Matrix, Vector, Vector >;
/// Construct a system solver.
/// \param[in] simulator The opm-models simulator object
/// \param[in] parameters Explicit parameters for solver setup, do not
/// read them from command line parameters.
ISTLSolverEbosWithGPU(const Simulator& simulator, const FlowLinearSolverParameters& parameters)
: ParentType(simulator)
{
bool have_gpu = true;
this->initialize(have_gpu);
}
/// Construct a system solver.
/// \param[in] simulator The opm-models simulator object
explicit ISTLSolverEbosWithGPU(const Simulator& simulator)
/// Construct a system solver.
/// \param[in] simulator The opm-models simulator object
/// \param[in] parameters Explicit parameters for solver setup, do not
/// read them from command line parameters.
ISTLSolverEbosWithGPU(const Simulator& simulator, const FlowLinearSolverParameters& parameters)
: ParentType(simulator)
{
}
{
bool have_gpu = true;
this->initialize(have_gpu);
}
void initialize()
/// Construct a system solver.
/// \param[in] simulator The opm-models simulator object
explicit ISTLSolverEbosWithGPU(const Simulator& simulator)
: ParentType(simulator)
{
}
void initialize()
{
OPM_TIMEBLOCK(initialize);
ParentType::initialize(false);
const bool on_io_rank = (this->simulator_.gridView().comm().rank() == 0);
{
OPM_TIMEBLOCK(initialize);
ParentType::initialize(false);
const bool on_io_rank = (this->simulator_.gridView().comm().rank() == 0);
{
std::string accelerator_mode = EWOMS_GET_PARAM(TypeTag, std::string, AcceleratorMode);
if ((this->simulator_.vanguard().grid().comm().size() > 1) && (accelerator_mode != "none")) {
if (on_io_rank) {
OpmLog::warning("Cannot use GPU with MPI, GPU are disabled");
}
accelerator_mode = "none";
std::string accelerator_mode = EWOMS_GET_PARAM(TypeTag, std::string, AcceleratorMode);
if ((this->simulator_.vanguard().grid().comm().size() > 1) && (accelerator_mode != "none")) {
if (on_io_rank) {
OpmLog::warning("Cannot use GPU with MPI, GPU are disabled");
}
const int platformID = EWOMS_GET_PARAM(TypeTag, int, OpenclPlatformId);
const int deviceID = EWOMS_GET_PARAM(TypeTag, int, BdaDeviceId);
const int maxit = EWOMS_GET_PARAM(TypeTag, int, LinearSolverMaxIter);
const double tolerance = EWOMS_GET_PARAM(TypeTag, double, LinearSolverReduction);
const bool opencl_ilu_parallel = EWOMS_GET_PARAM(TypeTag, bool, OpenclIluParallel);
const int linear_solver_verbosity = this->parameters_.linear_solver_verbosity_;
std::string linsolver = EWOMS_GET_PARAM(TypeTag, std::string, LinearSolver);
bdaBridge_ = std::make_unique<detail::BdaSolverInfo<Matrix,Vector>>(accelerator_mode,
linear_solver_verbosity,
maxit,
tolerance,
platformID,
deviceID,
opencl_ilu_parallel,
linsolver);
accelerator_mode = "none";
}
const int platformID = EWOMS_GET_PARAM(TypeTag, int, OpenclPlatformId);
const int deviceID = EWOMS_GET_PARAM(TypeTag, int, BdaDeviceId);
const int maxit = EWOMS_GET_PARAM(TypeTag, int, LinearSolverMaxIter);
const double tolerance = EWOMS_GET_PARAM(TypeTag, double, LinearSolverReduction);
const bool opencl_ilu_parallel = EWOMS_GET_PARAM(TypeTag, bool, OpenclIluParallel);
const int linear_solver_verbosity = this->parameters_.linear_solver_verbosity_;
std::string linsolver = EWOMS_GET_PARAM(TypeTag, std::string, LinearSolver);
bdaBridge_ = std::make_unique<detail::BdaSolverInfo<Matrix,Vector>>(accelerator_mode,
linear_solver_verbosity,
maxit,
tolerance,
platformID,
deviceID,
opencl_ilu_parallel,
linsolver);
}
}
void prepare(const Matrix& M, Vector& b)
{
OPM_TIMEBLOCK(prepare);
ParentType::prepare(M,b);
const bool firstcall = (this->matrix_ == nullptr);
// update matrix entries for solvers.
if (firstcall) {
// ebos will not change the matrix object. Hence simply store a pointer
// to the original one with a deleter that does nothing.
// Outch! We need to be able to scale the linear system! Hence const_cast
// setup sparsity pattern for jacobi matrix for preconditioner (only used for openclSolver)
void prepare(const Matrix& M, Vector& b)
{
OPM_TIMEBLOCK(prepare);
ParentType::prepare(M,b);
const bool firstcall = (this->matrix_ == nullptr);
// update matrix entries for solvers.
if (firstcall) {
// ebos will not change the matrix object. Hence simply store a pointer
// to the original one with a deleter that does nothing.
// Outch! We need to be able to scale the linear system! Hence const_cast
// setup sparsity pattern for jacobi matrix for preconditioner (only used for openclSolver)
#if HAVE_OPENCL
bdaBridge_->numJacobiBlocks_ = EWOMS_GET_PARAM(TypeTag, int, NumJacobiBlocks);
bdaBridge_->prepare(this->simulator_.vanguard().grid(),
this->simulator_.vanguard().cartesianIndexMapper(),
this->simulator_.vanguard().schedule().getWellsatEnd(),
this->simulator_.vanguard().cellPartition(),
this->getMatrix().nonzeroes(), this->useWellConn_);
bdaBridge_->numJacobiBlocks_ = EWOMS_GET_PARAM(TypeTag, int, NumJacobiBlocks);
bdaBridge_->prepare(this->simulator_.vanguard().grid(),
this->simulator_.vanguard().cartesianIndexMapper(),
this->simulator_.vanguard().schedule().getWellsatEnd(),
this->simulator_.vanguard().cellPartition(),
this->getMatrix().nonzeroes(), this->useWellConn_);
#endif
}
}
}
void setResidual(Vector& /* b */)
{
// rhs_ = &b; // Must be handled in prepare() instead.
}
void getResidual(Vector& b) const
{
b = *(this->rhs_);
}
void setMatrix(const SparseMatrixAdapter& /* M */)
{
// matrix_ = &M.istlMatrix(); // Must be handled in prepare() instead.
}
bool solve(Vector& x)
{
OPM_TIMEBLOCK(solve);
this->calls_ += 1;
// Write linear system if asked for.
const int verbosity = this->prm_.template get<int>("verbosity", 0);
const bool write_matrix = verbosity > 10;
if (write_matrix) {
Helper::writeSystem(this->simulator_, //simulator is only used to get names
this->getMatrix(),
*(this->rhs_),
this->comm_.get());
}
// Solve system.
Dune::InverseOperatorResult result;
void setResidual(Vector& /* b */)
{
// rhs_ = &b; // Must be handled in prepare() instead.
}
void getResidual(Vector& b) const
{
b = *(this->rhs_);
}
void setMatrix(const SparseMatrixAdapter& /* M */)
{
// matrix_ = &M.istlMatrix(); // Must be handled in prepare() instead.
}
bool solve(Vector& x)
{
OPM_TIMEBLOCK(solve);
this->calls_ += 1;
// Write linear system if asked for.
const int verbosity = this->prm_.template get<int>("verbosity", 0);
const bool write_matrix = verbosity > 10;
if (write_matrix) {
Helper::writeSystem(this->simulator_, //simulator is only used to get names
this->getMatrix(),
*(this->rhs_),
this->comm_.get());
}
// Solve system.
Dune::InverseOperatorResult result;
std::function<void(WellContributions&)> getContribs =
[this](WellContributions& w)
{
this->simulator_.problem().wellModel().getWellContributions(w);
};
if (!bdaBridge_->apply(*(this->rhs_), this->useWellConn_, getContribs,
this->simulator_.gridView().comm().rank(),
const_cast<Matrix&>(this->getMatrix()),
x, result))
std::function<void(WellContributions&)> getContribs =
[this](WellContributions& w)
{
if(bdaBridge_->gpuActive()){
// bda solve fails use istl solver setup need to be done since it is not setup in prepeare
ParentType::prepareFlexibleSolver();
}
assert(this->flexibleSolver_.solver_);
this->flexibleSolver_.solver_->apply(x, *(this->rhs_), result);
}
// Check convergence, iterations etc.
this->checkConvergence(result);
return this->converged_;
}
protected:
void prepareFlexibleSolver()
this->simulator_.problem().wellModel().getWellContributions(w);
};
if (!bdaBridge_->apply(*(this->rhs_), this->useWellConn_, getContribs,
this->simulator_.gridView().comm().rank(),
const_cast<Matrix&>(this->getMatrix()),
x, result))
{
if(bdaBridge_->gpuActive()){
// bda solve fails use istl solver setup need to be done since it is not setup in prepare
ParentType::prepareFlexibleSolver();
}
assert(this->flexibleSolver_.solver_);
this->flexibleSolver_.solver_->apply(x, *(this->rhs_), result);
}
std::unique_ptr<detail::BdaSolverInfo<Matrix, Vector>> bdaBridge_;
}; // end ISTLSolver
// Check convergence, iterations etc.
this->checkConvergence(result);
return this->converged_;
}
protected:
void prepareFlexibleSolver()
{
if(bdaBridge_->gpuActive()){
ParentType::prepareFlexibleSolver();
}
}
std::unique_ptr<detail::BdaSolverInfo<Matrix, Vector>> bdaBridge_;
}; // end ISTLSolver
} // namespace Opm
#endif
#endif
#endif // OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED

6
opm/simulators/wells/MultisegmentWellEquations.cpp
View File

@ -28,7 +28,10 @@
#include <opm/input/eclipse/Schedule/MSW/WellSegments.hpp>
#if COMPILE_BDA_BRIDGE
#include <opm/simulators/linalg/bda/WellContributions.hpp>
#endif
#include <opm/simulators/linalg/istlsparsematrixadapter.hh>
#include <opm/simulators/linalg/matrixblock.hh>
#include <opm/simulators/linalg/SmallDenseMatrixUtils.hpp>
@ -189,6 +192,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const
xw = mswellhelpers::applyUMFPack(*duneDSolver_, resWell);
}
#if COMPILE_BDA_BRIDGE
template<class Scalar, int numWellEq, int numEq>
void MultisegmentWellEquations<Scalar,numWellEq,numEq>::
extract(WellContributions& wellContribs) const
@ -255,7 +259,7 @@ extract(WellContributions& wellContribs) const
Drows,
Cvals);
}
#endif
template<class Scalar, int numWellEq, int numEq>
template<class SparseMatrixAdapter>

4
opm/simulators/wells/MultisegmentWellEquations.hpp
View File

@ -38,7 +38,9 @@ namespace Opm
template<class Scalar, int numWellEq, int numEq> class MultisegmentWellEquationAccess;
template<class Scalar> class MultisegmentWellGeneric;
#if COMPILE_BDA_BRIDGE
class WellContributions;
#endif
class WellInterfaceGeneric;
class WellState;
@ -98,8 +100,10 @@ public:
//! \details xw = inv(D)*(rw - C*x)
void recoverSolutionWell(const BVector& x, BVectorWell& xw) const;
#if COMPILE_BDA_BRIDGE
//! \brief Add the matrices of this well to the WellContributions object.
void extract(WellContributions& wellContribs) const;
#endif
//! \brief Add the matrices of this well to the sparse matrix adapter.
template<class SparseMatrixAdapter>

7
opm/simulators/wells/StandardWellEquations.cpp
View File

@ -24,9 +24,10 @@
#include <opm/common/TimingMacros.hpp>
#include <opm/simulators/wells/StandardWellEquations.hpp>
#if COMPILE_BDA_BRIDGE
#include <opm/simulators/linalg/bda/WellContributions.hpp>
#endif
#include <opm/simulators/linalg/istlsparsematrixadapter.hh>
#include <opm/simulators/linalg/matrixblock.hh>
#include <opm/simulators/linalg/SmallDenseMatrixUtils.hpp>
@ -187,6 +188,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const
invDuneD_.mv(resWell, xw);
}
#if COMPILE_BDA_BRIDGE
template<class Scalar, int numEq>
void StandardWellEquations<Scalar,numEq>::
extract(const int numStaticWellEq,
@ -240,6 +242,7 @@ extract(const int numStaticWellEq,
wellContribs.addMatrix(WellContributions::MatrixType::B,
colIndices.data(), nnzValues.data(), duneB_.nonzeroes());
}
#endif
template<class Scalar, int numEq>
template<class SparseMatrixAdapter>

4
opm/simulators/wells/StandardWellEquations.hpp
View File

@ -36,7 +36,9 @@ namespace Opm
class ParallelWellInfo;
template<class Scalar, int numEq> class StandardWellEquationAccess;
#if COMPILE_BDA_BRIDGE
class WellContributions;
#endif
class WellInterfaceGeneric;
class WellState;
@ -94,9 +96,11 @@ public:
//! \details xw = inv(D)*(rw - C*x)
void recoverSolutionWell(const BVector& x, BVectorWell& xw) const;
#if COMPILE_BDA_BRIDGE
//! \brief Add the matrices of this well to the WellContributions object.
void extract(const int numStaticWellEq,
WellContributions& wellContribs) const;
#endif
//! \brief Add the matrices of this well to the sparse matrix adapter.
template<class SparseMatrixAdapter>