OilfieldToolsNet/opm-simulators
4
0
Fork 0
mirror of https://github.com/OPM/opm-simulators.git synced 2024-07-04 11:33:06 -05:00
Code Issues Packages Projects Releases Wiki Activity

WellContribution: template Scalar type

Browse Source
This commit is contained in:
Arne Morten Kvarving 2024-04-15 16:57:54 +02:00
parent 216f0bea0d
commit ff39bc3b03
31 changed files with 137 additions and 98 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
opm/simulators/linalg/ISTLSolverBda.cpp
View File

@ -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<double>::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

8
opm/simulators/linalg/ISTLSolverBda.hpp
View File

@ -35,13 +35,13 @@ namespace Opm {
class Well;
template<class Matrix, class Vector, int block_size> class BdaBridge;
class WellContributions;
template<class Scalar> class WellContributions;
namespace detail {
template<class Matrix, class Vector>
struct BdaSolverInfo
{
using WellContribFunc = std::function<void(WellContributions&)>;
using WellContribFunc = std::function<void(WellContributions<double>&)>;
using Bridge = BdaBridge<Matrix,Vector,Matrix::block_type::rows>;
BdaSolverInfo(const std::string& accelerator_mode,
@ -249,8 +249,8 @@ public:
// Solve system.
Dune::InverseOperatorResult result;
std::function<void(WellContributions&)> getContribs =
[this](WellContributions& w)
std::function<void(WellContributions<double>&)> getContribs =
[this](WellContributions<double>& w)
{
this->simulator_.problem().wellModel().getWellContributions(w);
};

19
opm/simulators/linalg/bda/BdaBridge.cpp
View File

@ -201,12 +201,13 @@ void checkMemoryContiguous(const BridgeMatrix& mat) {
template <class BridgeMatrix, class BridgeVector, int block_size>
void BdaBridge<BridgeMatrix, BridgeVector, block_size>::solve_system(BridgeMatrix* bridgeMat,
BridgeMatrix* jacMat,
int numJacobiBlocks,
BridgeVector& b,
WellContributions& wellContribs,
InverseOperatorResult& res)
void BdaBridge<BridgeMatrix, BridgeVector, block_size>::
solve_system(BridgeMatrix* bridgeMat,
BridgeMatrix* jacMat,
int numJacobiBlocks,
BridgeVector& b,
WellContributions<double>& wellContribs,
InverseOperatorResult& res)
{
if (use_gpu) {
BdaResult result;
@ -302,8 +303,10 @@ void BdaBridge<BridgeMatrix, BridgeVector, block_size>::get_result([[maybe_unuse
}
template <class BridgeMatrix, class BridgeVector, int block_size>
void BdaBridge<BridgeMatrix, BridgeVector, block_size>::initWellContributions([[maybe_unused]] WellContributions& wellContribs,
[[maybe_unused]] unsigned N) {
void BdaBridge<BridgeMatrix, BridgeVector, block_size>::
initWellContributions([[maybe_unused]] WellContributions<double>& wellContribs,
[[maybe_unused]] unsigned N)
{
if(accelerator_mode.compare("opencl") == 0){
#if HAVE_OPENCL
const auto openclBackend = static_cast<const Opm::Accelerator::openclSolverBackend<block_size>*>(backend.get());

13
opm/simulators/linalg/bda/BdaBridge.hpp
View File

@ -27,7 +27,7 @@
namespace Opm
{
class WellContributions;
template<class Scalar> 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<Accelerator::BdaSolver<block_size> > backend;
std::unique_ptr<Accelerator::BdaSolver<block_size>> backend;
std::shared_ptr<Accelerator::BlockedMatrix<double>> matrix; // 'stores' matrix, actually points to h_rows, h_cols and the received BridgeMatrix for the nonzeroes
std::shared_ptr<Accelerator::BlockedMatrix<double>> jacMatrix; // 'stores' preconditioner matrix, actually points to h_rows, h_cols and the received BridgeMatrix for the nonzeroes
std::vector<int> 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<double>& 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<double>& wellContribs, unsigned N);
/// Return the selected accelerator mode, this is input via the command-line
std::string getAccleratorName(){

5
opm/simulators/linalg/bda/BdaSolver.hpp
View File

@ -29,7 +29,7 @@
namespace Opm {
class WellContributions;
template<class Scalar> class WellContributions;
namespace Accelerator {
enum class SolverStatus {
@ -86,7 +86,8 @@ namespace Accelerator {
virtual SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs, BdaResult& res) = 0;
WellContributions<double>& wellContribs,
BdaResult& res) = 0;
virtual void get_result(double *x) = 0;

59
opm/simulators/linalg/bda/WellContributions.cpp
View File

@ -39,10 +39,12 @@
namespace Opm {
WellContributions::~WellContributions() = default;
template<class Scalar>
WellContributions<Scalar>::~WellContributions() = default;
std::unique_ptr<WellContributions>
WellContributions::create(const std::string& accelerator_mode, bool useWellConn)
template<class Scalar>
std::unique_ptr<WellContributions<Scalar>>
WellContributions<Scalar>::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<class Scalar>
void WellContributions<Scalar>::
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<class Scalar>
void WellContributions<Scalar>::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<class Scalar>
void WellContributions<Scalar>::setVectorSize(unsigned N_)
{
N = N_;
}
void WellContributions::addNumBlocks(unsigned int numBlocks)
template<class Scalar>
void WellContributions<Scalar>::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<class Scalar>
void WellContributions<Scalar>::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<double>& Bvalues,
std::vector<unsigned int>& BcolIndices,
std::vector<unsigned int>& BrowPointers,
unsigned int DnumBlocks,
double* Dvalues,
UMFPackIndex* DcolPointers,
UMFPackIndex* DrowIndices,
std::vector<double>& Cvalues)
template<class Scalar>
void WellContributions<Scalar>::
addMultisegmentWellContribution(unsigned int dim_,
unsigned int dim_wells_,
unsigned int Mb,
std::vector<Scalar>& Bvalues,
std::vector<unsigned int>& BcolIndices,
std::vector<unsigned int>& BrowPointers,
unsigned int DnumBlocks,
Scalar* Dvalues,
UMFPackIndex* DcolPointers,
UMFPackIndex* DrowIndices,
std::vector<Scalar>& Cvalues)
{
assert(dim==dim_);
using MSW = MultisegmentWellContribution<double>;
using MSW = MultisegmentWellContribution<Scalar>;
multisegments.push_back(std::make_unique<MSW>(dim_,
dim_wells_,
Mb,
@ -172,4 +183,6 @@ void WellContributions::addMultisegmentWellContribution(unsigned int dim_,
++num_ms_wells;
}
template class WellContributions<double>;
} //namespace Opm

23
opm/simulators/linalg/bda/WellContributions.hpp
View File

@ -48,6 +48,7 @@ template<class Scalar> class MultisegmentWellContribution;
/// - get total size of all wellcontributions that must be stored here
/// - allocate memory
/// - copy data of wellcontributions
template<class Scalar>
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<unsigned int> val_pointers; // val_pointers[wellID] == index of first block for this well in Ccols and Bcols
std::vector<std::unique_ptr<MultisegmentWellContribution<double>>> multisegments;
std::vector<std::unique_ptr<MultisegmentWellContribution<Scalar>>> 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<double> &Bvalues, std::vector<unsigned int> &BcolIndices, std::vector<unsigned int> &BrowPointers,
unsigned int DnumBlocks, double *Dvalues,
UMFPackIndex *DcolPointers, UMFPackIndex *DrowIndices,
std::vector<double> &Cvalues);
std::vector<Scalar>& Bvalues,
std::vector<unsigned int>& BcolIndices,
std::vector<unsigned int>& BrowPointers,
unsigned int DnumBlocks,
Scalar* Dvalues,
UMFPackIndex* DcolPointers,
UMFPackIndex* DrowIndices,
std::vector<Scalar>& 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

2
opm/simulators/linalg/bda/amgclSolverBackend.cpp
View File

@ -408,7 +408,7 @@ SolverStatus amgclSolverBackend<block_size>::
solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
[[maybe_unused]] std::shared_ptr<BlockedMatrix<double>> jacMatrix,
[[maybe_unused]] WellContributions& wellContribs,
[[maybe_unused]] WellContributions<double>& wellContribs,
BdaResult& res)
{
if (initialized == false) {

2
opm/simulators/linalg/bda/amgclSolverBackend.hpp
View File

@ -143,7 +143,7 @@ public:
SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res) override;
/// Get result after linear solve, and peform postprocessing if necessary

2
opm/simulators/linalg/bda/cuda/cuWellContributions.hpp
View File

@ -28,7 +28,7 @@
namespace Opm
{
class WellContributionsCuda : public WellContributions
class WellContributionsCuda : public WellContributions<double>
{
public:
~WellContributionsCuda() override;

10
opm/simulators/linalg/bda/cuda/cusparseSolverBackend.cu
View File

@ -90,7 +90,9 @@ cusparseSolverBackend<block_size>::~cusparseSolverBackend() {
}
template <unsigned int block_size>
void cusparseSolverBackend<block_size>::gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) {
void cusparseSolverBackend<block_size>::
gpu_pbicgstab(WellContributions<double>& 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<block_size>::create_preconditioner() {
template <unsigned int block_size>
void cusparseSolverBackend<block_size>::solve_system(WellContributions& wellContribs, BdaResult &res) {
void cusparseSolverBackend<block_size>::
solve_system(WellContributions<double>& wellContribs, BdaResult& res)
{
// actually solve
gpu_pbicgstab(wellContribs, res);
cudaStreamSynchronize(stream);
@ -538,7 +542,7 @@ SolverStatus cusparseSolverBackend<block_size>::
solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res)
{
if (initialized == false) {

7
opm/simulators/linalg/bda/cuda/cusparseSolverBackend.hpp
View File

@ -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<double>& 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<double>& wellContribs, BdaResult &res);
public:
@ -142,7 +142,8 @@ public:
SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs, BdaResult& res) override;
WellContributions<double>& 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

11
opm/simulators/linalg/bda/opencl/openclSolverBackend.cpp
View File

@ -233,7 +233,9 @@ void openclSolverBackend<block_size>::setOpencl(std::shared_ptr<cl::Context>& co
template <unsigned int block_size>
void openclSolverBackend<block_size>::gpu_pbicgstab(WellContributions& wellContribs, BdaResult& res) {
void openclSolverBackend<block_size>::
gpu_pbicgstab(WellContributions<double>& wellContribs, BdaResult& res)
{
float it;
double rho, rhop, beta, alpha, omega, tmp1, tmp2;
double norm, norm_0;
@ -598,7 +600,9 @@ bool openclSolverBackend<block_size>::create_preconditioner() {
template <unsigned int block_size>
void openclSolverBackend<block_size>::solve_system(WellContributions &wellContribs, BdaResult &res) {
void openclSolverBackend<block_size>::
solve_system(WellContributions<double>& wellContribs, BdaResult& res)
{
Timer t;
// actually solve
@ -620,7 +624,6 @@ void openclSolverBackend<block_size>::solve_system(WellContributions &wellContri
out << "openclSolver::solve_system(): " << t.stop() << " s";
OpmLog::info(out.str());
}
} // end solve_system()
@ -644,7 +647,7 @@ SolverStatus openclSolverBackend<block_size>::
solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res)
{
if (initialized == false) {

6
opm/simulators/linalg/bda/opencl/openclSolverBackend.hpp
View File

@ -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<double>& 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<double>& wellContribs, BdaResult& res);
public:
std::shared_ptr<cl::Context> context;
@ -138,7 +138,7 @@ public:
SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res) override;
/// Solve scalar linear system, for example a coarse system of an AMG preconditioner

2
opm/simulators/linalg/bda/opencl/openclWellContributions.hpp
View File

@ -32,7 +32,7 @@
namespace Opm
{
class WellContributionsOCL : public WellContributions
class WellContributionsOCL : public WellContributions<double>
{
public:
void setOpenCLEnv(cl::Context *context_, cl::CommandQueue *queue_);

3
opm/simulators/linalg/bda/rocalutionSolverBackend.cpp
View File

@ -131,7 +131,6 @@ void rocalutionSolverBackend<block_size>::convert_matrix(BlockedMatrix<double>*
}
}
// copy result to host memory
// caller must be sure that x is a valid array
template <unsigned int block_size>
@ -152,7 +151,7 @@ SolverStatus rocalutionSolverBackend<block_size>::
solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
[[maybe_unused]] std::shared_ptr<BlockedMatrix<double>> jacMatrix,
[[maybe_unused]] WellContributions& wellContribs,
[[maybe_unused]] WellContributions<double>& wellContribs,
BdaResult& res)
{
if (initialized == false) {

2
opm/simulators/linalg/bda/rocalutionSolverBackend.hpp
View File

@ -94,7 +94,7 @@ public:
SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res) override;
/// Get result after linear solve, and peform postprocessing if necessary

11
opm/simulators/linalg/bda/rocsparseSolverBackend.cpp
View File

@ -149,8 +149,9 @@ rocsparseSolverBackend<block_size>::~rocsparseSolverBackend() {
template <unsigned int block_size>
void rocsparseSolverBackend<block_size>::gpu_pbicgstab([[maybe_unused]] WellContributions& wellContribs,
BdaResult& res)
void rocsparseSolverBackend<block_size>::
gpu_pbicgstab([[maybe_unused]] WellContributions<double>& wellContribs,
BdaResult& res)
{
float it = 0.5;
double rho, rhop, beta, alpha, nalpha, omega, nomega, tmp1, tmp2;
@ -602,7 +603,9 @@ bool rocsparseSolverBackend<block_size>::create_preconditioner() {
template <unsigned int block_size>
void rocsparseSolverBackend<block_size>::solve_system(WellContributions &wellContribs, BdaResult &res) {
void rocsparseSolverBackend<block_size>::
solve_system(WellContributions<double>& wellContribs, BdaResult& res)
{
Timer t;
// actually solve
@ -640,7 +643,7 @@ SolverStatus rocsparseSolverBackend<block_size>::
solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double *b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res)
{
if (initialized == false) {

6
opm/simulators/linalg/bda/rocsparseSolverBackend.hpp
View File

@ -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<double>& 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<double>& wellContribs, BdaResult& res);
public:
/// Construct a openclSolver
@ -141,7 +141,7 @@ public:
SolverStatus solve_system(std::shared_ptr<BlockedMatrix<double>> matrix,
double* b,
std::shared_ptr<BlockedMatrix<double>> jacMatrix,
WellContributions& wellContribs,
WellContributions<double>& wellContribs,
BdaResult& res) override;
/// Solve scalar linear system, for example a coarse system of an AMG preconditioner

2
opm/simulators/linalg/bda/rocsparseWellContributions.hpp
View File

@ -30,7 +30,7 @@
namespace Opm
{
class WellContributionsRocsparse : public WellContributions
class WellContributionsRocsparse : public WellContributions<double>
{
private:
hipStream_t stream;

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

@ -90,7 +90,7 @@ struct EnableTerminalOutput {
namespace Opm {
#if COMPILE_BDA_BRIDGE
class WellContributions;
template<class Scalar> 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<Scalar>& x) const;
#endif
// apply well model with scaling of alpha

2
opm/simulators/wells/BlackoilWellModel_impl.hpp
View File

@ -1569,7 +1569,7 @@ namespace Opm {
template<typename TypeTag>
void
BlackoilWellModel<TypeTag>::
getWellContributions(WellContributions& wellContribs) const
getWellContributions(WellContributions<Scalar>& wellContribs) const
{
// prepare for StandardWells
wellContribs.setBlockSize(StandardWell<TypeTag>::Indices::numEq, StandardWell<TypeTag>::numStaticWellEq);

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

@ -202,7 +202,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const
#if COMPILE_BDA_BRIDGE
template<class Scalar, int numWellEq, int numEq>
void MultisegmentWellEquations<Scalar,numWellEq,numEq>::
extract(WellContributions& wellContribs) const
extract(WellContributions<Scalar>& wellContribs) const
{
unsigned int Mb = duneB_.N(); // number of blockrows in duneB_, duneC_ and duneD_
unsigned int BnumBlocks = duneB_.nonzeroes();

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

@ -39,7 +39,7 @@ namespace Opm
template<class Scalar, int numWellEq, int numEq> class MultisegmentWellEquationAccess;
template<class Scalar> class MultisegmentWellGeneric;
#if COMPILE_BDA_BRIDGE
class WellContributions;
template<class Scalar> class WellContributions;
#endif
template<class Scalar> class WellInterfaceGeneric;
template<class Scalar> 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<Scalar>& wellContribs) const;
#endif
//! \brief Add the matrices of this well to the sparse matrix adapter.

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

@ -198,7 +198,7 @@ recoverSolutionWell(const BVector& x, BVectorWell& xw) const
template<class Scalar, int numEq>
void StandardWellEquations<Scalar,numEq>::
extract(const int numStaticWellEq,
WellContributions& wellContribs) const
WellContributions<Scalar>& wellContribs) const
{
std::vector<int> colIndices;
std::vector<Scalar> nnzValues;
@ -216,7 +216,7 @@ extract(const int numStaticWellEq,
}
}
}
wellContribs.addMatrix(WellContributions::MatrixType::C,
wellContribs.addMatrix(WellContributions<Scalar>::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<Scalar>::MatrixType::D,
colIndices.data(), nnzValues.data(), 1);
// duneB
@ -245,7 +245,7 @@ extract(const int numStaticWellEq,
}
}
}
wellContribs.addMatrix(WellContributions::MatrixType::B,
wellContribs.addMatrix(WellContributions<Scalar>::MatrixType::B,
colIndices.data(), nnzValues.data(), duneB_.nonzeroes());
}
#endif

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

@ -37,7 +37,7 @@ namespace Opm
template<class Scalar> class ParallelWellInfo;
template<class Scalar, int numEq> class StandardWellEquationAccess;
#if COMPILE_BDA_BRIDGE
class WellContributions;
template<class Scalar> class WellContributions;
#endif
template<class Scalar> class WellInterfaceGeneric;
template<class Scalar> 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<Scalar>& wellContribs) const;
#endif
//! \brief Add the matrices of this well to the sparse matrix adapter.

2
opm/simulators/wells/StandardWellEval.hpp
View File

@ -38,7 +38,7 @@ class ConvergenceReport;
class DeferredLogger;
class Schedule;
class SummaryState;
class WellContributions;
template<class Scalar> class WellContributions;
template<class FluidSystem, class Indices> class WellInterfaceIndices;
template<class Scalar> class WellState;

4
tests/test_cusparseSolver.cpp
View File

@ -123,7 +123,7 @@ testCusparseSolver(Opm::BdaBridge<Matrix<bz>, Vector<bz>, bz>& bridge, Matrix<bz
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("cusparse", false);
auto wellContribs = Opm::WellContributions<double>::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<Matrix<bz>, Vector<bz>, bz>& bridge, Mat
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("cusparse", false);
auto wellContribs = Opm::WellContributions<double>::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

4
tests/test_openclSolver.cpp
View File

@ -120,7 +120,7 @@ testOpenclSolver(Opm::BdaBridge<Matrix<bz>, Vector<bz>, bz>& bridge, Matrix<bz>&
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("opencl", false);
auto wellContribs = Opm::WellContributions<double>::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<Matrix<bz>, Vector<bz>, bz>& bridge, Matri
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("opencl", false);
auto wellContribs = Opm::WellContributions<double>::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

2
tests/test_rocalutionSolver.cpp
View File

@ -96,7 +96,7 @@ testRocalutionSolver(const boost::property_tree::ptree& prm, Matrix<bz>& matrix,
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create(accelerator_mode, true);
auto wellContribs = Opm::WellContributions<double>::create(accelerator_mode, true);
std::unique_ptr<Opm::BdaBridge<Matrix<bz>, Vector<bz>, bz> > bridge;
try {
bridge = std::make_unique<Opm::BdaBridge<Matrix<bz>, Vector<bz>, bz> >(accelerator_mode,

4
tests/test_rocsparseSolver.cpp
View File

@ -127,7 +127,7 @@ testRocsparseSolver(std::unique_ptr<Opm::BdaBridge<Matrix<bz>, Vector<bz>, bz> >
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("rocsparse", true);
auto wellContribs = Opm::WellContributions<double>::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<Opm::BdaBridge<Matrix<bz>, Vector<bz>,
{
Dune::InverseOperatorResult result;
Vector<bz> x(rhs.size());
auto wellContribs = Opm::WellContributions::create("rocsparse", true);
auto wellContribs = Opm::WellContributions<double>::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