From ff39bc3b0384588c0c2baca6881ad6907a91773d Mon Sep 17 00:00:00 2001 From: Arne Morten Kvarving Date: Mon, 15 Apr 2024 16:57:54 +0200 Subject: [PATCH] WellContribution: template Scalar type --- opm/simulators/linalg/ISTLSolverBda.cpp | 2 +- opm/simulators/linalg/ISTLSolverBda.hpp | 8 +-- opm/simulators/linalg/bda/BdaBridge.cpp | 19 +++--- opm/simulators/linalg/bda/BdaBridge.hpp | 13 ++-- opm/simulators/linalg/bda/BdaSolver.hpp | 5 +- .../linalg/bda/WellContributions.cpp | 59 +++++++++++-------- .../linalg/bda/WellContributions.hpp | 23 +++++--- .../linalg/bda/amgclSolverBackend.cpp | 2 +- .../linalg/bda/amgclSolverBackend.hpp | 2 +- .../linalg/bda/cuda/cuWellContributions.hpp | 2 +- .../linalg/bda/cuda/cusparseSolverBackend.cu | 10 +++- .../linalg/bda/cuda/cusparseSolverBackend.hpp | 7 ++- .../linalg/bda/opencl/openclSolverBackend.cpp | 11 ++-- .../linalg/bda/opencl/openclSolverBackend.hpp | 6 +- .../bda/opencl/openclWellContributions.hpp | 2 +- .../linalg/bda/rocalutionSolverBackend.cpp | 3 +- .../linalg/bda/rocalutionSolverBackend.hpp | 2 +- .../linalg/bda/rocsparseSolverBackend.cpp | 11 ++-- .../linalg/bda/rocsparseSolverBackend.hpp | 6 +- .../linalg/bda/rocsparseWellContributions.hpp | 2 +- opm/simulators/wells/BlackoilWellModel.hpp | 4 +- .../wells/BlackoilWellModel_impl.hpp | 2 +- .../wells/MultisegmentWellEquations.cpp | 2 +- .../wells/MultisegmentWellEquations.hpp | 4 +- .../wells/StandardWellEquations.cpp | 8 +-- .../wells/StandardWellEquations.hpp | 4 +- opm/simulators/wells/StandardWellEval.hpp | 2 +- tests/test_cusparseSolver.cpp | 4 +- tests/test_openclSolver.cpp | 4 +- tests/test_rocalutionSolver.cpp | 2 +- tests/test_rocsparseSolver.cpp | 4 +- 31 files changed, 137 insertions(+), 98 deletions(-) diff --git a/opm/simulators/linalg/ISTLSolverBda.cpp b/opm/simulators/linalg/ISTLSolverBda.cpp index 422905990..940c7ee11 100644 --- a/opm/simulators/linalg/ISTLSolverBda.cpp +++ b/opm/simulators/linalg/ISTLSolverBda.cpp @@ -104,7 +104,7 @@ apply(Vector& rhs, { bool use_gpu = bridge_->getUseGpu(); if (use_gpu) { - auto wellContribs = WellContributions::create(accelerator_mode_, useWellConn); + auto wellContribs = WellContributions::create(accelerator_mode_, useWellConn); bridge_->initWellContributions(*wellContribs, x.N() * x[0].N()); // the WellContributions can only be applied separately with CUDA, OpenCL or rocsparse, not with amgcl or rocalution diff --git a/opm/simulators/linalg/ISTLSolverBda.hpp b/opm/simulators/linalg/ISTLSolverBda.hpp index ca64d5494..98a03b58c 100644 --- a/opm/simulators/linalg/ISTLSolverBda.hpp +++ b/opm/simulators/linalg/ISTLSolverBda.hpp @@ -35,13 +35,13 @@ namespace Opm { class Well; template class BdaBridge; -class WellContributions; +template class WellContributions; namespace detail { template struct BdaSolverInfo { - using WellContribFunc = std::function; + using WellContribFunc = std::function&)>; using Bridge = BdaBridge; BdaSolverInfo(const std::string& accelerator_mode, @@ -249,8 +249,8 @@ public: // Solve system. Dune::InverseOperatorResult result; - std::function getContribs = - [this](WellContributions& w) + std::function&)> getContribs = + [this](WellContributions& w) { this->simulator_.problem().wellModel().getWellContributions(w); }; diff --git a/opm/simulators/linalg/bda/BdaBridge.cpp b/opm/simulators/linalg/bda/BdaBridge.cpp index b806402f8..c9f6254be 100644 --- a/opm/simulators/linalg/bda/BdaBridge.cpp +++ b/opm/simulators/linalg/bda/BdaBridge.cpp @@ -201,12 +201,13 @@ void checkMemoryContiguous(const BridgeMatrix& mat) { template -void BdaBridge::solve_system(BridgeMatrix* bridgeMat, - BridgeMatrix* jacMat, - int numJacobiBlocks, - BridgeVector& b, - WellContributions& wellContribs, - InverseOperatorResult& res) +void BdaBridge:: +solve_system(BridgeMatrix* bridgeMat, + BridgeMatrix* jacMat, + int numJacobiBlocks, + BridgeVector& b, + WellContributions& wellContribs, + InverseOperatorResult& res) { if (use_gpu) { BdaResult result; @@ -302,8 +303,10 @@ void BdaBridge::get_result([[maybe_unuse } template -void BdaBridge::initWellContributions([[maybe_unused]] WellContributions& wellContribs, - [[maybe_unused]] unsigned N) { +void BdaBridge:: +initWellContributions([[maybe_unused]] WellContributions& wellContribs, + [[maybe_unused]] unsigned N) +{ if(accelerator_mode.compare("opencl") == 0){ #if HAVE_OPENCL const auto openclBackend = static_cast*>(backend.get()); diff --git a/opm/simulators/linalg/bda/BdaBridge.hpp b/opm/simulators/linalg/bda/BdaBridge.hpp index df896d869..95d62e468 100644 --- a/opm/simulators/linalg/bda/BdaBridge.hpp +++ b/opm/simulators/linalg/bda/BdaBridge.hpp @@ -27,7 +27,7 @@ namespace Opm { -class WellContributions; +template class WellContributions; typedef Dune::InverseOperatorResult InverseOperatorResult; @@ -39,7 +39,7 @@ private: int verbosity = 0; bool use_gpu = false; std::string accelerator_mode; - std::unique_ptr > backend; + std::unique_ptr> backend; std::shared_ptr> matrix; // 'stores' matrix, actually points to h_rows, h_cols and the received BridgeMatrix for the nonzeroes std::shared_ptr> jacMatrix; // 'stores' preconditioner matrix, actually points to h_rows, h_cols and the received BridgeMatrix for the nonzeroes std::vector h_rows, h_cols; // store the sparsity pattern of the matrix @@ -69,7 +69,12 @@ public: /// \param[in] b vector b, should be of type Dune::BlockVector /// \param[in] wellContribs contains all WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] result summary of solver result - void solve_system(BridgeMatrix *bridgeMat, BridgeMatrix *jacMat, int numJacobiBlocks, BridgeVector &b, WellContributions& wellContribs, InverseOperatorResult &result); + void solve_system(BridgeMatrix* bridgeMat, + BridgeMatrix* jacMat, + int numJacobiBlocks, + BridgeVector& b, + WellContributions& wellContribs, + InverseOperatorResult &result); /// Get the resulting x vector /// \param[inout] x vector x, should be of type Dune::BlockVector @@ -91,7 +96,7 @@ public: /// those must be set before calling BlackOilWellModel::getWellContributions() in ISTL /// \param[in] wellContribs container to hold all WellContributions /// \param[in] N number of rows in scalar vector that wellContribs will be applied on - void initWellContributions(WellContributions& wellContribs, unsigned N); + void initWellContributions(WellContributions& wellContribs, unsigned N); /// Return the selected accelerator mode, this is input via the command-line std::string getAccleratorName(){ diff --git a/opm/simulators/linalg/bda/BdaSolver.hpp b/opm/simulators/linalg/bda/BdaSolver.hpp index e700d0451..0d9e67f42 100644 --- a/opm/simulators/linalg/bda/BdaSolver.hpp +++ b/opm/simulators/linalg/bda/BdaSolver.hpp @@ -29,7 +29,7 @@ namespace Opm { -class WellContributions; +template class WellContributions; namespace Accelerator { enum class SolverStatus { @@ -86,7 +86,8 @@ namespace Accelerator { virtual SolverStatus solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, BdaResult& res) = 0; + WellContributions& wellContribs, + BdaResult& res) = 0; virtual void get_result(double *x) = 0; diff --git a/opm/simulators/linalg/bda/WellContributions.cpp b/opm/simulators/linalg/bda/WellContributions.cpp index 66436d2ae..c23e7b256 100644 --- a/opm/simulators/linalg/bda/WellContributions.cpp +++ b/opm/simulators/linalg/bda/WellContributions.cpp @@ -39,10 +39,12 @@ namespace Opm { -WellContributions::~WellContributions() = default; +template +WellContributions::~WellContributions() = default; -std::unique_ptr -WellContributions::create(const std::string& accelerator_mode, bool useWellConn) +template +std::unique_ptr> +WellContributions::create(const std::string& accelerator_mode, bool useWellConn) { if(accelerator_mode.compare("cusparse") == 0){ #if HAVE_CUDA @@ -86,10 +88,12 @@ WellContributions::create(const std::string& accelerator_mode, bool useWellConn) } } -void WellContributions::addMatrix([[maybe_unused]] MatrixType type, - [[maybe_unused]] int* colIndices, - [[maybe_unused]] double* values, - [[maybe_unused]] unsigned int val_size) +template +void WellContributions:: +addMatrix([[maybe_unused]] MatrixType type, + [[maybe_unused]] int* colIndices, + [[maybe_unused]] Scalar* values, + [[maybe_unused]] unsigned int val_size) { #if !HAVE_CUDA && !HAVE_OPENCL OPM_THROW(std::logic_error, "Error cannot add StandardWell matrix on GPU because neither CUDA nor OpenCL were found by cmake"); @@ -107,7 +111,8 @@ void WellContributions::addMatrix([[maybe_unused]] MatrixType type, } } -void WellContributions::setBlockSize(unsigned int dim_, unsigned int dim_wells_) +template +void WellContributions::setBlockSize(unsigned int dim_, unsigned int dim_wells_) { dim = dim_; dim_wells = dim_wells_; @@ -121,11 +126,14 @@ void WellContributions::setBlockSize(unsigned int dim_, unsigned int dim_wells_) } } -void WellContributions::setVectorSize(unsigned N_) { +template +void WellContributions::setVectorSize(unsigned N_) +{ N = N_; } -void WellContributions::addNumBlocks(unsigned int numBlocks) +template +void WellContributions::addNumBlocks(unsigned int numBlocks) { if (allocated) { OPM_THROW(std::logic_error, "Error cannot add more sizes after allocated in WellContributions"); @@ -134,7 +142,8 @@ void WellContributions::addNumBlocks(unsigned int numBlocks) num_std_wells++; } -void WellContributions::alloc() +template +void WellContributions::alloc() { if (num_std_wells > 0) { val_pointers.resize(num_std_wells+1); @@ -144,20 +153,22 @@ void WellContributions::alloc() } } -void WellContributions::addMultisegmentWellContribution(unsigned int dim_, - unsigned int dim_wells_, - unsigned int Mb, - std::vector& Bvalues, - std::vector& BcolIndices, - std::vector& BrowPointers, - unsigned int DnumBlocks, - double* Dvalues, - UMFPackIndex* DcolPointers, - UMFPackIndex* DrowIndices, - std::vector& Cvalues) +template +void WellContributions:: +addMultisegmentWellContribution(unsigned int dim_, + unsigned int dim_wells_, + unsigned int Mb, + std::vector& Bvalues, + std::vector& BcolIndices, + std::vector& BrowPointers, + unsigned int DnumBlocks, + Scalar* Dvalues, + UMFPackIndex* DcolPointers, + UMFPackIndex* DrowIndices, + std::vector& Cvalues) { assert(dim==dim_); - using MSW = MultisegmentWellContribution; + using MSW = MultisegmentWellContribution; multisegments.push_back(std::make_unique(dim_, dim_wells_, Mb, @@ -172,4 +183,6 @@ void WellContributions::addMultisegmentWellContribution(unsigned int dim_, ++num_ms_wells; } +template class WellContributions; + } //namespace Opm diff --git a/opm/simulators/linalg/bda/WellContributions.hpp b/opm/simulators/linalg/bda/WellContributions.hpp index 282e37b5a..eb908b748 100644 --- a/opm/simulators/linalg/bda/WellContributions.hpp +++ b/opm/simulators/linalg/bda/WellContributions.hpp @@ -48,6 +48,7 @@ template class MultisegmentWellContribution; /// - get total size of all wellcontributions that must be stored here /// - allocate memory /// - copy data of wellcontributions +template class WellContributions { public: @@ -74,7 +75,7 @@ protected: unsigned int num_std_wells_so_far = 0; // keep track of where next data is written std::vector val_pointers; // val_pointers[wellID] == index of first block for this well in Ccols and Bcols - std::vector>> multisegments; + std::vector>> multisegments; public: unsigned int getNumWells(){ @@ -105,7 +106,7 @@ public: /// \param[in] colIndices columnindices of blocks in C or B, ignored for D /// \param[in] values array of nonzeroes /// \param[in] val_size number of blocks in C or B, ignored for D - void addMatrix(MatrixType type, int *colIndices, double *values, unsigned int val_size); + void addMatrix(MatrixType type, int* colIndices, Scalar* values, unsigned int val_size); /// Add a MultisegmentWellContribution, actually creates an object on heap that is destroyed in the destructor /// Matrices C and B are passed in Blocked CSR, matrix D in CSC @@ -120,19 +121,25 @@ public: /// \param[in] DcolPointers columnpointers of matrix D /// \param[in] DrowIndices rowindices of matrix D /// \param[in] Cvalues nonzero values of matrix C - void addMultisegmentWellContribution(unsigned int dim, unsigned int dim_wells, + void addMultisegmentWellContribution(unsigned int dim, + unsigned int dim_wells, unsigned int Mb, - std::vector &Bvalues, std::vector &BcolIndices, std::vector &BrowPointers, - unsigned int DnumBlocks, double *Dvalues, - UMFPackIndex *DcolPointers, UMFPackIndex *DrowIndices, - std::vector &Cvalues); + std::vector& Bvalues, + std::vector& BcolIndices, + std::vector& BrowPointers, + unsigned int DnumBlocks, + Scalar* Dvalues, + UMFPackIndex* DcolPointers, + UMFPackIndex* DrowIndices, + std::vector& Cvalues); protected: //! \brief API specific allocation. virtual void APIalloc() {} /// Api specific upload of matrix. - virtual void APIaddMatrix(MatrixType, int*, double*, unsigned int) {} + virtual void APIaddMatrix(MatrixType, int*, Scalar*, unsigned int) {} }; + } //namespace Opm #endif diff --git a/opm/simulators/linalg/bda/amgclSolverBackend.cpp b/opm/simulators/linalg/bda/amgclSolverBackend.cpp index 8fac42bdb..8d4c86fb6 100644 --- a/opm/simulators/linalg/bda/amgclSolverBackend.cpp +++ b/opm/simulators/linalg/bda/amgclSolverBackend.cpp @@ -408,7 +408,7 @@ SolverStatus amgclSolverBackend:: solve_system(std::shared_ptr> matrix, double *b, [[maybe_unused]] std::shared_ptr> jacMatrix, - [[maybe_unused]] WellContributions& wellContribs, + [[maybe_unused]] WellContributions& wellContribs, BdaResult& res) { if (initialized == false) { diff --git a/opm/simulators/linalg/bda/amgclSolverBackend.hpp b/opm/simulators/linalg/bda/amgclSolverBackend.hpp index f7c01955b..0908d1e9c 100644 --- a/opm/simulators/linalg/bda/amgclSolverBackend.hpp +++ b/opm/simulators/linalg/bda/amgclSolverBackend.hpp @@ -143,7 +143,7 @@ public: SolverStatus solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) override; /// Get result after linear solve, and peform postprocessing if necessary diff --git a/opm/simulators/linalg/bda/cuda/cuWellContributions.hpp b/opm/simulators/linalg/bda/cuda/cuWellContributions.hpp index 892a2411b..2e5e84547 100644 --- a/opm/simulators/linalg/bda/cuda/cuWellContributions.hpp +++ b/opm/simulators/linalg/bda/cuda/cuWellContributions.hpp @@ -28,7 +28,7 @@ namespace Opm { -class WellContributionsCuda : public WellContributions +class WellContributionsCuda : public WellContributions { public: ~WellContributionsCuda() override; diff --git a/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu b/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu index 4e2b14e15..3d9d8b1e3 100644 --- a/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu +++ b/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu @@ -90,7 +90,9 @@ cusparseSolverBackend::~cusparseSolverBackend() { } template -void cusparseSolverBackend::gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) { +void cusparseSolverBackend:: +gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) +{ Timer t_total, t_prec(false), t_spmv(false), t_well(false), t_rest(false); int n = N; double rho = 1.0, rhop; @@ -509,7 +511,9 @@ bool cusparseSolverBackend::create_preconditioner() { template -void cusparseSolverBackend::solve_system(WellContributions& wellContribs, BdaResult &res) { +void cusparseSolverBackend:: +solve_system(WellContributions& wellContribs, BdaResult& res) +{ // actually solve gpu_pbicgstab(wellContribs, res); cudaStreamSynchronize(stream); @@ -538,7 +542,7 @@ SolverStatus cusparseSolverBackend:: solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) { if (initialized == false) { diff --git a/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.hpp b/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.hpp index a4a51b9f6..cfb2ccab8 100644 --- a/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.hpp +++ b/opm/simulators/linalg/bda/cuda/cusparseSolverBackend.hpp @@ -77,7 +77,7 @@ private: /// Solve linear system using ilu0-bicgstab /// \param[in] wellContribs contains all WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] res summary of solver result - void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); + void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); /// Initialize GPU and allocate memory /// \param[in] matrix matrix for spmv @@ -117,7 +117,7 @@ private: /// Solve linear system /// \param[in] wellContribs contains all WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] res summary of solver result - void solve_system(WellContributions& wellContribs, BdaResult &res); + void solve_system(WellContributions& wellContribs, BdaResult &res); public: @@ -142,7 +142,8 @@ public: SolverStatus solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, BdaResult& res) override; + WellContributions& wellContribs, + BdaResult& res) override; /// Get resulting vector x after linear solve, also includes post processing if necessary /// \param[inout] x resulting x vector, caller must guarantee that x points to a valid array diff --git a/opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp b/opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp index 1e34ee83f..eaf0c2d72 100644 --- a/opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp +++ b/opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp @@ -233,7 +233,9 @@ void openclSolverBackend::setOpencl(std::shared_ptr& co template -void openclSolverBackend::gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) { +void openclSolverBackend:: +gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) +{ float it; double rho, rhop, beta, alpha, omega, tmp1, tmp2; double norm, norm_0; @@ -598,7 +600,9 @@ bool openclSolverBackend::create_preconditioner() { template -void openclSolverBackend::solve_system(WellContributions &wellContribs, BdaResult &res) { +void openclSolverBackend:: +solve_system(WellContributions& wellContribs, BdaResult& res) +{ Timer t; // actually solve @@ -620,7 +624,6 @@ void openclSolverBackend::solve_system(WellContributions &wellContri out << "openclSolver::solve_system(): " << t.stop() << " s"; OpmLog::info(out.str()); } - } // end solve_system() @@ -644,7 +647,7 @@ SolverStatus openclSolverBackend:: solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) { if (initialized == false) { diff --git a/opm/simulators/linalg/bda/opencl/openclSolverBackend.hpp b/opm/simulators/linalg/bda/opencl/openclSolverBackend.hpp index 9e63d6098..f0d0d73a0 100644 --- a/opm/simulators/linalg/bda/opencl/openclSolverBackend.hpp +++ b/opm/simulators/linalg/bda/opencl/openclSolverBackend.hpp @@ -76,7 +76,7 @@ private: /// Solve linear system using ilu0-bicgstab /// \param[in] wellContribs WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] res summary of solver result - void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); + void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); /// Initialize GPU and allocate memory /// \param[in] matrix matrix A @@ -107,7 +107,7 @@ private: /// \param[in] wellContribs WellContributions, to apply them separately, instead of adding them to matrix A /// could be empty /// \param[inout] res summary of solver result - void solve_system(WellContributions &wellContribs, BdaResult &res); + void solve_system(WellContributions& wellContribs, BdaResult& res); public: std::shared_ptr context; @@ -138,7 +138,7 @@ public: SolverStatus solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) override; /// Solve scalar linear system, for example a coarse system of an AMG preconditioner diff --git a/opm/simulators/linalg/bda/opencl/openclWellContributions.hpp b/opm/simulators/linalg/bda/opencl/openclWellContributions.hpp index 7c877ccc6..60e27fb4a 100644 --- a/opm/simulators/linalg/bda/opencl/openclWellContributions.hpp +++ b/opm/simulators/linalg/bda/opencl/openclWellContributions.hpp @@ -32,7 +32,7 @@ namespace Opm { -class WellContributionsOCL : public WellContributions +class WellContributionsOCL : public WellContributions { public: void setOpenCLEnv(cl::Context *context_, cl::CommandQueue *queue_); diff --git a/opm/simulators/linalg/bda/rocalutionSolverBackend.cpp b/opm/simulators/linalg/bda/rocalutionSolverBackend.cpp index d4a579bf8..5cc4b4f5f 100644 --- a/opm/simulators/linalg/bda/rocalutionSolverBackend.cpp +++ b/opm/simulators/linalg/bda/rocalutionSolverBackend.cpp @@ -131,7 +131,6 @@ void rocalutionSolverBackend::convert_matrix(BlockedMatrix* } } - // copy result to host memory // caller must be sure that x is a valid array template @@ -152,7 +151,7 @@ SolverStatus rocalutionSolverBackend:: solve_system(std::shared_ptr> matrix, double *b, [[maybe_unused]] std::shared_ptr> jacMatrix, - [[maybe_unused]] WellContributions& wellContribs, + [[maybe_unused]] WellContributions& wellContribs, BdaResult& res) { if (initialized == false) { diff --git a/opm/simulators/linalg/bda/rocalutionSolverBackend.hpp b/opm/simulators/linalg/bda/rocalutionSolverBackend.hpp index 8e853edfa..a7450eeca 100644 --- a/opm/simulators/linalg/bda/rocalutionSolverBackend.hpp +++ b/opm/simulators/linalg/bda/rocalutionSolverBackend.hpp @@ -94,7 +94,7 @@ public: SolverStatus solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) override; /// Get result after linear solve, and peform postprocessing if necessary diff --git a/opm/simulators/linalg/bda/rocsparseSolverBackend.cpp b/opm/simulators/linalg/bda/rocsparseSolverBackend.cpp index 844a36860..74ede44e1 100644 --- a/opm/simulators/linalg/bda/rocsparseSolverBackend.cpp +++ b/opm/simulators/linalg/bda/rocsparseSolverBackend.cpp @@ -149,8 +149,9 @@ rocsparseSolverBackend::~rocsparseSolverBackend() { template -void rocsparseSolverBackend::gpu_pbicgstab([[maybe_unused]] WellContributions& wellContribs, - BdaResult& res) +void rocsparseSolverBackend:: +gpu_pbicgstab([[maybe_unused]] WellContributions& wellContribs, + BdaResult& res) { float it = 0.5; double rho, rhop, beta, alpha, nalpha, omega, nomega, tmp1, tmp2; @@ -602,7 +603,9 @@ bool rocsparseSolverBackend::create_preconditioner() { template -void rocsparseSolverBackend::solve_system(WellContributions &wellContribs, BdaResult &res) { +void rocsparseSolverBackend:: +solve_system(WellContributions& wellContribs, BdaResult& res) +{ Timer t; // actually solve @@ -640,7 +643,7 @@ SolverStatus rocsparseSolverBackend:: solve_system(std::shared_ptr> matrix, double *b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) { if (initialized == false) { diff --git a/opm/simulators/linalg/bda/rocsparseSolverBackend.hpp b/opm/simulators/linalg/bda/rocsparseSolverBackend.hpp index 319b26d3e..2301415f2 100644 --- a/opm/simulators/linalg/bda/rocsparseSolverBackend.hpp +++ b/opm/simulators/linalg/bda/rocsparseSolverBackend.hpp @@ -87,7 +87,7 @@ private: /// Solve linear system using ilu0-bicgstab /// \param[in] wellContribs WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] res summary of solver result - void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); + void gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res); /// Initialize GPU and allocate memory /// \param[in] matrix matrix A @@ -114,7 +114,7 @@ private: /// Solve linear system /// \param[in] wellContribs WellContributions, to apply them separately, instead of adding them to matrix A /// \param[inout] res summary of solver result - void solve_system(WellContributions &wellContribs, BdaResult &res); + void solve_system(WellContributions& wellContribs, BdaResult& res); public: /// Construct a openclSolver @@ -141,7 +141,7 @@ public: SolverStatus solve_system(std::shared_ptr> matrix, double* b, std::shared_ptr> jacMatrix, - WellContributions& wellContribs, + WellContributions& wellContribs, BdaResult& res) override; /// Solve scalar linear system, for example a coarse system of an AMG preconditioner diff --git a/opm/simulators/linalg/bda/rocsparseWellContributions.hpp b/opm/simulators/linalg/bda/rocsparseWellContributions.hpp index bfa7e888a..1a80bf7ad 100644 --- a/opm/simulators/linalg/bda/rocsparseWellContributions.hpp +++ b/opm/simulators/linalg/bda/rocsparseWellContributions.hpp @@ -30,7 +30,7 @@ namespace Opm { -class WellContributionsRocsparse : public WellContributions +class WellContributionsRocsparse : public WellContributions { private: hipStream_t stream; diff --git a/opm/simulators/wells/BlackoilWellModel.hpp b/opm/simulators/wells/BlackoilWellModel.hpp index 0d4c07487..eed486c42 100644 --- a/opm/simulators/wells/BlackoilWellModel.hpp +++ b/opm/simulators/wells/BlackoilWellModel.hpp @@ -90,7 +90,7 @@ struct EnableTerminalOutput { namespace Opm { #if COMPILE_BDA_BRIDGE -class WellContributions; +template class WellContributions; #endif /// Class for handling the blackoil well model. @@ -287,7 +287,7 @@ class WellContributions; #if COMPILE_BDA_BRIDGE // accumulate the contributions of all Wells in the WellContributions object - void getWellContributions(WellContributions& x) const; + void getWellContributions(WellContributions& x) const; #endif // apply well model with scaling of alpha diff --git a/opm/simulators/wells/BlackoilWellModel_impl.hpp b/opm/simulators/wells/BlackoilWellModel_impl.hpp index f2cf467db..2374abffc 100644 --- a/opm/simulators/wells/BlackoilWellModel_impl.hpp +++ b/opm/simulators/wells/BlackoilWellModel_impl.hpp @@ -1569,7 +1569,7 @@ namespace Opm { template void BlackoilWellModel:: - getWellContributions(WellContributions& wellContribs) const + getWellContributions(WellContributions& wellContribs) const { // prepare for StandardWells wellContribs.setBlockSize(StandardWell::Indices::numEq, StandardWell::numStaticWellEq); diff --git a/opm/simulators/wells/MultisegmentWellEquations.cpp b/opm/simulators/wells/MultisegmentWellEquations.cpp index 0f86ff51b..39ebe5c79 100644 --- a/opm/simulators/wells/MultisegmentWellEquations.cpp +++ b/opm/simulators/wells/MultisegmentWellEquations.cpp @@ -202,7 +202,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const #if COMPILE_BDA_BRIDGE template void MultisegmentWellEquations:: -extract(WellContributions& wellContribs) const +extract(WellContributions& wellContribs) const { unsigned int Mb = duneB_.N(); // number of blockrows in duneB_, duneC_ and duneD_ unsigned int BnumBlocks = duneB_.nonzeroes(); diff --git a/opm/simulators/wells/MultisegmentWellEquations.hpp b/opm/simulators/wells/MultisegmentWellEquations.hpp index 58289e48e..f204c0699 100644 --- a/opm/simulators/wells/MultisegmentWellEquations.hpp +++ b/opm/simulators/wells/MultisegmentWellEquations.hpp @@ -39,7 +39,7 @@ namespace Opm template class MultisegmentWellEquationAccess; template class MultisegmentWellGeneric; #if COMPILE_BDA_BRIDGE -class WellContributions; +template class WellContributions; #endif template class WellInterfaceGeneric; template class WellState; @@ -105,7 +105,7 @@ public: #if COMPILE_BDA_BRIDGE //! \brief Add the matrices of this well to the WellContributions object. - void extract(WellContributions& wellContribs) const; + void extract(WellContributions& wellContribs) const; #endif //! \brief Add the matrices of this well to the sparse matrix adapter. diff --git a/opm/simulators/wells/StandardWellEquations.cpp b/opm/simulators/wells/StandardWellEquations.cpp index bcce8c414..9ead9bf35 100644 --- a/opm/simulators/wells/StandardWellEquations.cpp +++ b/opm/simulators/wells/StandardWellEquations.cpp @@ -198,7 +198,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const template void StandardWellEquations:: extract(const int numStaticWellEq, - WellContributions& wellContribs) const + WellContributions& wellContribs) const { std::vector colIndices; std::vector nnzValues; @@ -216,7 +216,7 @@ extract(const int numStaticWellEq, } } } - wellContribs.addMatrix(WellContributions::MatrixType::C, + wellContribs.addMatrix(WellContributions::MatrixType::C, colIndices.data(), nnzValues.data(), duneC_.nonzeroes()); // invDuneD @@ -229,7 +229,7 @@ extract(const int numStaticWellEq, nnzValues.emplace_back(invDuneD_[0][0][i][j]); } } - wellContribs.addMatrix(WellContributions::MatrixType::D, + wellContribs.addMatrix(WellContributions::MatrixType::D, colIndices.data(), nnzValues.data(), 1); // duneB @@ -245,7 +245,7 @@ extract(const int numStaticWellEq, } } } - wellContribs.addMatrix(WellContributions::MatrixType::B, + wellContribs.addMatrix(WellContributions::MatrixType::B, colIndices.data(), nnzValues.data(), duneB_.nonzeroes()); } #endif diff --git a/opm/simulators/wells/StandardWellEquations.hpp b/opm/simulators/wells/StandardWellEquations.hpp index ae7c26bbd..736d7ded1 100644 --- a/opm/simulators/wells/StandardWellEquations.hpp +++ b/opm/simulators/wells/StandardWellEquations.hpp @@ -37,7 +37,7 @@ namespace Opm template class ParallelWellInfo; template class StandardWellEquationAccess; #if COMPILE_BDA_BRIDGE -class WellContributions; +template class WellContributions; #endif template class WellInterfaceGeneric; template class WellState; @@ -102,7 +102,7 @@ public: #if COMPILE_BDA_BRIDGE //! \brief Add the matrices of this well to the WellContributions object. void extract(const int numStaticWellEq, - WellContributions& wellContribs) const; + WellContributions& wellContribs) const; #endif //! \brief Add the matrices of this well to the sparse matrix adapter. diff --git a/opm/simulators/wells/StandardWellEval.hpp b/opm/simulators/wells/StandardWellEval.hpp index a5d62ffdb..c9fb6f5c6 100644 --- a/opm/simulators/wells/StandardWellEval.hpp +++ b/opm/simulators/wells/StandardWellEval.hpp @@ -38,7 +38,7 @@ class ConvergenceReport; class DeferredLogger; class Schedule; class SummaryState; -class WellContributions; +template class WellContributions; template class WellInterfaceIndices; template class WellState; diff --git a/tests/test_cusparseSolver.cpp b/tests/test_cusparseSolver.cpp index a6724df89..573556165 100644 --- a/tests/test_cusparseSolver.cpp +++ b/tests/test_cusparseSolver.cpp @@ -123,7 +123,7 @@ testCusparseSolver(Opm::BdaBridge, Vector, bz>& bridge, Matrix x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("cusparse", false); + auto wellContribs = Opm::WellContributions::create("cusparse", false); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory bridge.solve_system(&mat2, &mat2, /*numJacobiBlocks=*/0, rhs, *wellContribs, result); @@ -138,7 +138,7 @@ testCusparseSolverJacobi(Opm::BdaBridge, Vector, bz>& bridge, Mat { Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("cusparse", false); + auto wellContribs = Opm::WellContributions::create("cusparse", false); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory auto mat3 = matrix; // another deep copy, to make sure Jacobi matrix memory is different diff --git a/tests/test_openclSolver.cpp b/tests/test_openclSolver.cpp index ff97ab4ad..b55efb078 100644 --- a/tests/test_openclSolver.cpp +++ b/tests/test_openclSolver.cpp @@ -120,7 +120,7 @@ testOpenclSolver(Opm::BdaBridge, Vector, bz>& bridge, Matrix& { Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("opencl", false); + auto wellContribs = Opm::WellContributions::create("opencl", false); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory bridge.solve_system(&mat2, &mat2, /*numJacobiBlocks=*/0, rhs, *wellContribs, result); @@ -135,7 +135,7 @@ testOpenclSolverJacobi(Opm::BdaBridge, Vector, bz>& bridge, Matri { Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("opencl", false); + auto wellContribs = Opm::WellContributions::create("opencl", false); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory auto mat3 = matrix; // another deep copy, to make sure Jacobi matrix memory is different diff --git a/tests/test_rocalutionSolver.cpp b/tests/test_rocalutionSolver.cpp index de7cb46b7..6489c4f61 100644 --- a/tests/test_rocalutionSolver.cpp +++ b/tests/test_rocalutionSolver.cpp @@ -96,7 +96,7 @@ testRocalutionSolver(const boost::property_tree::ptree& prm, Matrix& matrix, Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create(accelerator_mode, true); + auto wellContribs = Opm::WellContributions::create(accelerator_mode, true); std::unique_ptr, Vector, bz> > bridge; try { bridge = std::make_unique, Vector, bz> >(accelerator_mode, diff --git a/tests/test_rocsparseSolver.cpp b/tests/test_rocsparseSolver.cpp index 05bb185c9..67000d2ce 100644 --- a/tests/test_rocsparseSolver.cpp +++ b/tests/test_rocsparseSolver.cpp @@ -127,7 +127,7 @@ testRocsparseSolver(std::unique_ptr, Vector, bz> > { Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("rocsparse", true); + auto wellContribs = Opm::WellContributions::create("rocsparse", true); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory bridge->solve_system(&mat2, &mat2, /*numJacobiBlocks=*/0, rhs, *wellContribs, result); @@ -142,7 +142,7 @@ testRocsparseSolverJacobi(std::unique_ptr, Vector, { Dune::InverseOperatorResult result; Vector x(rhs.size()); - auto wellContribs = Opm::WellContributions::create("rocsparse", true); + auto wellContribs = Opm::WellContributions::create("rocsparse", true); auto mat2 = matrix; // deep copy to make sure nnz values are in contiguous memory // matrix created by readMatrixMarket() did not have contiguous memory auto mat3 = matrix; // another deep copy, to make sure Jacobi matrix memory is different