mirror of
https://github.com/OPM/opm-simulators.git
synced 2025-01-20 01:32:57 -06:00
BdaBridge and WellContributions are only included and compiled when CUDA is found
This commit is contained in:
T.D. (Tongdong) Qiu
parent
d29b6a6e5e
commit
fdcf46792a
@ -29,7 +29,6 @@ list (APPEND MAIN_SOURCE_FILES
|
||||
opm/simulators/flow/MissingFeatures.cpp
|
||||
opm/simulators/linalg/ExtractParallelGridInformationToISTL.cpp
|
||||
opm/simulators/linalg/setupPropertyTree.cpp
|
||||
opm/simulators/linalg/bda/BdaBridge.cpp
|
||||
opm/simulators/timestepping/TimeStepControl.cpp
|
||||
opm/simulators/timestepping/AdaptiveSimulatorTimer.cpp
|
||||
opm/simulators/timestepping/SimulatorTimer.cpp
|
||||
@ -46,6 +45,7 @@ list (APPEND MAIN_SOURCE_FILES
|
||||
if(CUDA_FOUND)
|
||||
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/cusparseSolverBackend.cu)
|
||||
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/WellContributions.cu)
|
||||
list (APPEND MAIN_SOURCE_FILES opm/simulators/linalg/bda/BdaBridge.cpp)
|
||||
endif()
|
||||
|
||||
# originally generated with the command:
|
||||
|
||||
@ -45,7 +45,9 @@
|
||||
|
||||
#include <opm/common/utility/platform_dependent/reenable_warnings.h>
|
||||
|
||||
#if HAVE_CUDA
|
||||
#include <opm/simulators/linalg/bda/BdaBridge.hpp>
|
||||
#endif
|
||||
|
||||
BEGIN_PROPERTIES
|
||||
|
||||
|
||||
@ -35,7 +35,6 @@ typedef Dune::InverseOperatorResult InverseOperatorResult;
|
||||
namespace Opm
|
||||
{
|
||||
|
||||
#if HAVE_CUDA
|
||||
BdaBridge::BdaBridge(bool use_gpu_, int linear_solver_verbosity, int maxit, double tolerance)
|
||||
: use_gpu(use_gpu_)
|
||||
{
|
||||
@ -43,15 +42,9 @@ BdaBridge::BdaBridge(bool use_gpu_, int linear_solver_verbosity, int maxit, doub
|
||||
backend.reset(new cusparseSolverBackend(linear_solver_verbosity, maxit, tolerance));
|
||||
}
|
||||
}
|
||||
#else
|
||||
BdaBridge::BdaBridge(bool use_gpu_ OPM_UNUSED, int linear_solver_verbosity OPM_UNUSED, int maxit OPM_UNUSED, double tolerance OPM_UNUSED)
|
||||
{
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
|
||||
#if HAVE_CUDA
|
||||
template <class BridgeMatrix>
|
||||
int checkZeroDiagonal(BridgeMatrix& mat) {
|
||||
static std::vector<typename BridgeMatrix::size_type> diag_indices; // contains offsets of the diagonal nnzs
|
||||
@ -118,13 +111,11 @@ void getSparsityPattern(BridgeMatrix& mat, std::vector<int> &h_rows, std::vector
|
||||
}
|
||||
} // end getSparsityPattern()
|
||||
|
||||
#endif
|
||||
|
||||
template <class BridgeMatrix, class BridgeVector>
|
||||
void BdaBridge::solve_system(BridgeMatrix *mat OPM_UNUSED, BridgeVector &b OPM_UNUSED, WellContributions& wellContribs OPM_UNUSED, InverseOperatorResult &res OPM_UNUSED)
|
||||
{
|
||||
|
||||
#if HAVE_CUDA
|
||||
if (use_gpu) {
|
||||
BdaResult result;
|
||||
result.converged = false;
|
||||
@ -192,17 +183,14 @@ void BdaBridge::solve_system(BridgeMatrix *mat OPM_UNUSED, BridgeVector &b OPM_U
|
||||
}else{
|
||||
res.converged = false;
|
||||
}
|
||||
#endif // HAVE_CUDA
|
||||
}
|
||||
|
||||
|
||||
template <class BridgeVector>
|
||||
void BdaBridge::get_result(BridgeVector &x OPM_UNUSED) {
|
||||
#if HAVE_CUDA
|
||||
if (use_gpu) {
|
||||
backend->post_process(static_cast<double*>(&(x[0][0])));
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
template void BdaBridge::solve_system<
|
||||
|
||||
@ -21,6 +21,11 @@
|
||||
#define BDABRIDGE_HEADER_INCLUDED
|
||||
|
||||
#include <config.h>
|
||||
|
||||
#if ! HAVE_CUDA
|
||||
#error "This file should only be included if CUDA is found"
|
||||
#endif
|
||||
|
||||
#include "dune/istl/solver.hh" // for struct InverseOperatorResult
|
||||
|
||||
#include "dune/istl/bcrsmatrix.hh"
|
||||
@ -28,9 +33,7 @@
|
||||
|
||||
#include <opm/simulators/linalg/bda/WellContributions.hpp>
|
||||
|
||||
#if HAVE_CUDA
|
||||
#include <opm/simulators/linalg/bda/cusparseSolverBackend.hpp>
|
||||
#endif
|
||||
|
||||
namespace Opm
|
||||
{
|
||||
@ -42,10 +45,8 @@ typedef Dune::InverseOperatorResult InverseOperatorResult;
|
||||
class BdaBridge
|
||||
{
|
||||
private:
|
||||
#if HAVE_CUDA
|
||||
std::unique_ptr<cusparseSolverBackend> backend;
|
||||
bool use_gpu;
|
||||
#endif
|
||||
|
||||
public:
|
||||
/// Construct a BdaBridge
|
||||
|
||||
@ -27,13 +27,14 @@
|
||||
#include <cuda_runtime.h>
|
||||
#endif
|
||||
|
||||
#include <opm/common/OpmLog/OpmLog.hpp>
|
||||
#include <opm/common/ErrorMacros.hpp>
|
||||
#include "opm/simulators/linalg/bda/WellContributions.hpp"
|
||||
|
||||
namespace Opm
|
||||
{
|
||||
|
||||
// apply WellContributions using y -= C^T * (D^-1 * (B * x))
|
||||
#if HAVE_CUDA
|
||||
__global__ void apply_well_contributions(
|
||||
const double * __restrict__ Cnnzs,
|
||||
const double * __restrict__ Dnnzs,
|
||||
@ -127,7 +128,6 @@ namespace Opm
|
||||
}
|
||||
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
void WellContributions::alloc(){
|
||||
|
||||
@ -17,16 +17,18 @@
|
||||
along with OPM. If not, see <http://www.gnu.org/licenses/>.
|
||||
*/
|
||||
|
||||
#ifndef WellContributions_H
|
||||
#define WellContributions_H
|
||||
#ifndef WELLCONTRIBUTIONS_HEADER_INCLUDED
|
||||
#define WELLCONTRIBUTIONS_HEADER_INCLUDED
|
||||
|
||||
#include <config.h>
|
||||
|
||||
#if ! HAVE_CUDA
|
||||
#error "This file should only be included if CUDA is found"
|
||||
#endif
|
||||
|
||||
#include <vector>
|
||||
|
||||
#if HAVE_CUDA
|
||||
#include <cuda_runtime.h>
|
||||
#endif
|
||||
|
||||
#include <config.h>
|
||||
|
||||
namespace Opm
|
||||
{
|
||||
@ -46,7 +48,6 @@ namespace Opm
|
||||
unsigned int *val_pointers = nullptr; // val_pointers[wellID] == index of first block for this well in Ccols and Bcols
|
||||
bool allocated = false;
|
||||
|
||||
#if HAVE_CUDA
|
||||
double *d_Cnnzs = nullptr;
|
||||
double *d_Dnnzs = nullptr;
|
||||
double *d_Bnnzs = nullptr;
|
||||
@ -56,23 +57,10 @@ namespace Opm
|
||||
double *d_z2 = nullptr;
|
||||
unsigned int *d_val_pointers = nullptr;
|
||||
cudaStream_t stream;
|
||||
#endif
|
||||
|
||||
double *Cnnzs = nullptr;
|
||||
double *Dnnzs = nullptr;
|
||||
double *Bnnzs = nullptr;
|
||||
int *Ccols = nullptr;
|
||||
int *Dcols = nullptr;
|
||||
int *Bcols = nullptr;
|
||||
double *z1 = nullptr; // z1 = B * x
|
||||
double *z2 = nullptr; // z2 = D^-1 * B * x
|
||||
|
||||
public:
|
||||
#if HAVE_CUDA
|
||||
/// Set a cudaStream to be used
|
||||
/// \param[in] stream the cudaStream that is used to launch the kernel in
|
||||
void setCudaStream(cudaStream_t stream);
|
||||
#endif
|
||||
|
||||
/// Create a new WellContributions, implementation is empty
|
||||
WellContributions(){};
|
||||
@ -93,7 +81,7 @@ namespace Opm
|
||||
void addMatrix(int idx, int *colIndices, double *values, unsigned int val_size);
|
||||
|
||||
/// Return the number of wells added to this object
|
||||
unsigned int get_num_wells(){
|
||||
unsigned int getNumWells(){
|
||||
return num_wells;
|
||||
}
|
||||
|
||||
|
||||
@ -76,7 +76,7 @@ namespace Opm
|
||||
|
||||
t_total1 = second();
|
||||
|
||||
if(wellContribs.get_num_wells() > 0){
|
||||
if(wellContribs.getNumWells() > 0){
|
||||
wellContribs.setCudaStream(stream);
|
||||
}
|
||||
|
||||
@ -120,7 +120,7 @@ namespace Opm
|
||||
&one, descr_M, d_bVals, d_bRows, d_bCols, block_size, d_pw, &zero, d_v);
|
||||
|
||||
// apply wellContributions
|
||||
if(wellContribs.get_num_wells() > 0){
|
||||
if(wellContribs.getNumWells() > 0){
|
||||
wellContribs.apply(d_pw, d_v);
|
||||
}
|
||||
|
||||
@ -151,7 +151,7 @@ namespace Opm
|
||||
d_bVals, d_bRows, d_bCols, block_size, d_s, &zero, d_t);
|
||||
|
||||
// apply wellContributions
|
||||
if(wellContribs.get_num_wells() > 0){
|
||||
if(wellContribs.getNumWells() > 0){
|
||||
wellContribs.apply(d_s, d_t);
|
||||
}
|
||||
|
||||
|
||||
@ -192,8 +192,10 @@ namespace Opm {
|
||||
// subtract B*inv(D)*C * x from A*x
|
||||
void apply(const BVector& x, BVector& Ax) const;
|
||||
|
||||
#if HAVE_CUDA
|
||||
// accumulate the contributions of all Wells in the WellContributions object
|
||||
void getWellContributions(WellContributions& x) const;
|
||||
#endif
|
||||
|
||||
// apply well model with scaling of alpha
|
||||
void applyScaleAdd(const Scalar alpha, const BVector& x, BVector& Ax) const;
|
||||
|
||||
@ -918,7 +918,7 @@ namespace Opm {
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
#if HAVE_CUDA
|
||||
template<typename TypeTag>
|
||||
void
|
||||
BlackoilWellModel<TypeTag>::
|
||||
@ -938,7 +938,7 @@ namespace Opm {
|
||||
derived->addWellContribution(wellContribs);
|
||||
}
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
// Ax = Ax - alpha * C D^-1 B x
|
||||
template<typename TypeTag>
|
||||
|
||||
@ -184,11 +184,13 @@ namespace Opm
|
||||
/// r = r - C D^-1 Rw
|
||||
virtual void apply(BVector& r) const override;
|
||||
|
||||
#if HAVE_CUDA
|
||||
/// add the contribution (C, D^-1, B matrices) of this Well to the WellContributions object
|
||||
void addWellContribution(WellContributions& wellContribs) const;
|
||||
|
||||
/// get the sizes of the C, D^-1 and B matrices, used to allocate memory in a WellContributions object
|
||||
void getWellSizes(unsigned int& _nnzs, unsigned int& _numEq, unsigned int& _numWellEq) const;
|
||||
#endif
|
||||
|
||||
/// using the solution x to recover the solution xw for wells and applying
|
||||
/// xw to update Well State
|
||||
|
||||
@ -2735,7 +2735,7 @@ namespace Opm
|
||||
duneC_.mmtv(invDrw_, r);
|
||||
}
|
||||
|
||||
|
||||
#if HAVE_CUDA
|
||||
template<typename TypeTag>
|
||||
void
|
||||
StandardWell<TypeTag>::
|
||||
@ -2795,7 +2795,7 @@ namespace Opm
|
||||
_numEq = numEq;
|
||||
_numWellEq = numStaticWellEq;
|
||||
}
|
||||
|
||||
#endif
|
||||
|
||||
|
||||
template<typename TypeTag>
|
||||
|
||||
Loading…
Reference in New Issue
Block a user