diff --git a/CMakeLists.txt b/CMakeLists.txt index 4750f7645..0dcb4c65f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) diff --git a/CMakeLists_files.cmake b/CMakeLists_files.cmake index 31b261fa2..1d3461d3e 100644 --- a/CMakeLists_files.cmake +++ b/CMakeLists_files.cmake @@ -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) diff --git a/opm/simulators/flow/BlackoilModelEbosNldd.hpp b/opm/simulators/flow/BlackoilModelEbosNldd.hpp index 27e2382fb..e715febd6 100644 --- a/opm/simulators/flow/BlackoilModelEbosNldd.hpp +++ b/opm/simulators/flow/BlackoilModelEbosNldd.hpp @@ -34,7 +34,12 @@ #include #include + +#if COMPILE_BDA_BRIDGE #include +#else +#include +#endif #include #include diff --git a/opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp b/opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp index 9df56dc2c..dc0a1f149 100644 --- a/opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp +++ b/opm/simulators/linalg/ISTLSolverEbosWithGPU.hpp @@ -21,14 +21,15 @@ #ifndef OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED #define OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED + #include -#if COMPILE_BDA_BRIDGE -namespace Opm -{ + +namespace Opm { + template class BdaBridge; class WellContributions; -namespace detail -{ +namespace detail { + template 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 ISTLSolverEbosWithGPU : public ISTLSolverEbos - { - protected: - using ParentType = ISTLSolverEbos; - using GridView = GetPropType; - using Scalar = GetPropType; - using SparseMatrixAdapter = GetPropType; - using Vector = GetPropType; - using Indices = GetPropType; - using WellModel = GetPropType; - using Simulator = GetPropType; - using Matrix = typename SparseMatrixAdapter::IstlMatrix; - using ThreadManager = GetPropType; - using ElementContext = GetPropType; - using AbstractSolverType = Dune::InverseOperator; - using AbstractOperatorType = Dune::AssembledLinearOperator; - using AbstractPreconditionerType = Dune::PreconditionerWithUpdate; - using WellModelOperator = WellModelAsLinearOperator; - using ElementMapper = GetPropType; - constexpr static std::size_t pressureIndex = GetPropType::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 ISTLSolverEbosWithGPU : public ISTLSolverEbos +{ +protected: + using ParentType = ISTLSolverEbos; + using GridView = GetPropType; + using Scalar = GetPropType; + using SparseMatrixAdapter = GetPropType; + using Vector = GetPropType; + using Indices = GetPropType; + using WellModel = GetPropType; + using Simulator = GetPropType; + using Matrix = typename SparseMatrixAdapter::IstlMatrix; + using ThreadManager = GetPropType; + using ElementContext = GetPropType; + using AbstractSolverType = Dune::InverseOperator; + using AbstractOperatorType = Dune::AssembledLinearOperator; + using AbstractPreconditionerType = Dune::PreconditionerWithUpdate; + using WellModelOperator = WellModelAsLinearOperator; + using ElementMapper = GetPropType; + constexpr static std::size_t pressureIndex = GetPropType::pressureSwitchIdx; #if HAVE_MPI - using CommunicationType = Dune::OwnerOverlapCopyCommunication; + using CommunicationType = Dune::OwnerOverlapCopyCommunication; #else - using CommunicationType = Dune::CollectiveCommunication; + using CommunicationType = Dune::CollectiveCommunication; #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>(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>(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("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("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 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(this->getMatrix()), - x, result)) + std::function 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(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> bdaBridge_; - }; // end ISTLSolver + + // Check convergence, iterations etc. + this->checkConvergence(result); + + return this->converged_; + } + +protected: + void prepareFlexibleSolver() + { + if(bdaBridge_->gpuActive()){ + ParentType::prepareFlexibleSolver(); + } + } + std::unique_ptr> bdaBridge_; +}; // end ISTLSolver } // namespace Opm -#endif -#endif + +#endif // OPM_ISTLSOLVER_EBOS_WITH_GPU_HEADER_INCLUDED diff --git a/opm/simulators/wells/MultisegmentWellEquations.cpp b/opm/simulators/wells/MultisegmentWellEquations.cpp index 4f08fb7d8..e84c8b43a 100644 --- a/opm/simulators/wells/MultisegmentWellEquations.cpp +++ b/opm/simulators/wells/MultisegmentWellEquations.cpp @@ -28,7 +28,10 @@ #include +#if COMPILE_BDA_BRIDGE #include +#endif + #include #include #include @@ -189,6 +192,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const xw = mswellhelpers::applyUMFPack(*duneDSolver_, resWell); } +#if COMPILE_BDA_BRIDGE template void MultisegmentWellEquations:: extract(WellContributions& wellContribs) const @@ -255,7 +259,7 @@ extract(WellContributions& wellContribs) const Drows, Cvals); } - +#endif template template diff --git a/opm/simulators/wells/MultisegmentWellEquations.hpp b/opm/simulators/wells/MultisegmentWellEquations.hpp index 4fcb76300..cef5d128d 100644 --- a/opm/simulators/wells/MultisegmentWellEquations.hpp +++ b/opm/simulators/wells/MultisegmentWellEquations.hpp @@ -38,7 +38,9 @@ namespace Opm template class MultisegmentWellEquationAccess; template 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 diff --git a/opm/simulators/wells/StandardWellEquations.cpp b/opm/simulators/wells/StandardWellEquations.cpp index f26883a00..9fc8f9733 100644 --- a/opm/simulators/wells/StandardWellEquations.cpp +++ b/opm/simulators/wells/StandardWellEquations.cpp @@ -24,9 +24,10 @@ #include #include - - +#if COMPILE_BDA_BRIDGE #include +#endif + #include #include #include @@ -187,6 +188,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const invDuneD_.mv(resWell, xw); } +#if COMPILE_BDA_BRIDGE template void StandardWellEquations:: 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 template diff --git a/opm/simulators/wells/StandardWellEquations.hpp b/opm/simulators/wells/StandardWellEquations.hpp index 503efe5f8..8ded42a72 100644 --- a/opm/simulators/wells/StandardWellEquations.hpp +++ b/opm/simulators/wells/StandardWellEquations.hpp @@ -36,7 +36,9 @@ namespace Opm class ParallelWellInfo; template 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