Merge pull request #3690 from akva2/wellcontrib_cleanups

Some simple cleanup in BDA WellContributions
This commit is contained in:
Markus Blatt
2021-11-15 10:18:12 +01:00
committed by GitHub
7 changed files with 41 additions and 31 deletions
+2 -1
View File
@@ -27,7 +27,6 @@
#include <opm/simulators/linalg/bda/BdaSolver.hpp>
#include <opm/simulators/linalg/bda/ILUReorder.hpp>
#include <opm/simulators/linalg/bda/WellContributions.hpp>
#if HAVE_FPGA
@@ -37,6 +36,8 @@
namespace Opm
{
class WellContributions;
typedef Dune::InverseOperatorResult InverseOperatorResult;
using Opm::Accelerator::ILUReorder;
+6 -7
View File
@@ -22,15 +22,14 @@
#include <opm/simulators/linalg/bda/BdaResult.hpp>
#include <opm/simulators/linalg/bda/WellContributions.hpp>
namespace Opm
{
namespace Accelerator
{
#include <string>
using Opm::WellContributions;
namespace Opm {
class WellContributions;
namespace Accelerator {
enum class SolverStatus {
BDA_SOLVER_SUCCESS,
BDA_SOLVER_ANALYSIS_FAILED,
@@ -57,7 +56,7 @@ namespace Accelerator
int maxit = 200;
double tolerance = 1e-2;
std::string bitstream = "";
std::string bitstream;
int N; // number of rows
int Nb; // number of blocked rows (Nb*block_size == N)
+25 -18
View File
@@ -54,10 +54,6 @@ WellContributions::WellContributions(std::string accelerator_mode, bool useWellC
WellContributions::~WellContributions()
{
// delete MultisegmentWellContributions
for (auto ms: multisegments) {
delete ms;
}
multisegments.clear();
#if HAVE_CUDA
@@ -66,10 +62,6 @@ WellContributions::~WellContributions()
}
#endif
if(num_std_wells > 0){
delete[] val_pointers;
}
#if HAVE_OPENCL
if(opencl_gpu){
if(num_ms_wells > 0){
@@ -121,7 +113,7 @@ void WellContributions::apply_mswells(cl::Buffer d_x, cl::Buffer d_y){
events.clear();
// actually apply MultisegmentWells
for(Opm::MultisegmentWellContribution *well: multisegments){
for (auto& well : multisegments) {
well->setReordering(h_toOrder, reorder);
well->apply(h_x, h_y);
}
@@ -185,7 +177,7 @@ void WellContributions::addMatrix([[maybe_unused]] MatrixType type, [[maybe_unus
if (num_std_wells_so_far == num_std_wells - 1) {
val_pointers[num_std_wells] = num_blocks;
events.resize(1);
queue->enqueueWriteBuffer(*d_val_pointers_ocl, CL_FALSE, 0, sizeof(unsigned int) * (num_std_wells + 1), val_pointers, nullptr, &events[0]);
queue->enqueueWriteBuffer(*d_val_pointers_ocl, CL_FALSE, 0, sizeof(unsigned int) * (num_std_wells + 1), val_pointers.data(), nullptr, &events[0]);
events[0].wait();
events.clear();
}
@@ -231,7 +223,7 @@ void WellContributions::addNumBlocks(unsigned int numBlocks)
void WellContributions::alloc()
{
if (num_std_wells > 0) {
val_pointers = new unsigned int[num_std_wells + 1];
val_pointers.resize(num_std_wells+1);
#if HAVE_CUDA
if(cuda_gpu){
@@ -253,15 +245,30 @@ 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)
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)
{
assert(dim==dim_);
MultisegmentWellContribution *well = new MultisegmentWellContribution(dim_, dim_wells_, Mb, Bvalues, BcolIndices, BrowPointers, DnumBlocks, Dvalues, DcolPointers, DrowIndices, Cvalues);
multisegments.emplace_back(well);
multisegments.push_back(std::make_unique<MultisegmentWellContribution>(dim_,
dim_wells_,
Mb,
Bvalues,
BcolIndices,
BrowPointers,
DnumBlocks,
Dvalues,
DcolPointers,
DrowIndices,
Cvalues));
++num_ms_wells;
}
@@ -177,7 +177,7 @@ void WellContributions::apply(double *d_x, double *d_y)
cudaStreamSynchronize(stream);
// actually apply MultisegmentWells
for (MultisegmentWellContribution *well : multisegments) {
for (auto& well : multisegments) {
well->apply(h_x, h_y);
}
@@ -210,7 +210,7 @@ void WellContributions::addMatrixGpu(MatrixType type, int *colIndices, double *v
val_pointers[num_std_wells_so_far] = num_blocks_so_far;
if (num_std_wells_so_far == num_std_wells - 1) {
val_pointers[num_std_wells] = num_blocks;
cudaMemcpy(d_val_pointers, val_pointers, sizeof(unsigned int) * (num_std_wells + 1), cudaMemcpyHostToDevice);
cudaMemcpy(d_val_pointers, val_pointers.data(), sizeof(unsigned int) * (num_std_wells + 1), cudaMemcpyHostToDevice);
}
break;
default:
@@ -222,7 +222,7 @@ void WellContributions::addMatrixGpu(MatrixType type, int *colIndices, double *v
void WellContributions::setCudaStream(cudaStream_t stream_)
{
this->stream = stream_;
for (MultisegmentWellContribution *well : multisegments) {
for (auto& well : multisegments) {
well->setCudaStream(stream_);
}
}
@@ -29,6 +29,7 @@
#include <opm/simulators/linalg/bda/openclKernels.hpp>
#endif
#include <memory>
#include <vector>
#include <opm/simulators/linalg/bda/MultisegmentWellContribution.hpp>
@@ -85,11 +86,11 @@ private:
unsigned int num_ms_wells = 0; // number of MultisegmentWells in this object, must equal multisegments.size()
unsigned int num_blocks_so_far = 0; // keep track of where next data is written
unsigned int num_std_wells_so_far = 0; // keep track of where next data is written
unsigned int *val_pointers = nullptr; // val_pointers[wellID] == index of first block for this well in Ccols and Bcols
std::vector<unsigned int> val_pointers; // val_pointers[wellID] == index of first block for this well in Ccols and Bcols
double *h_x = nullptr;
double *h_y = nullptr;
std::vector<MultisegmentWellContribution*> multisegments;
std::vector<std::unique_ptr<MultisegmentWellContribution>> multisegments;
#if HAVE_OPENCL
cl::Context *context;
+1
View File
@@ -30,6 +30,7 @@
BOOST_VERSION / 100 % 1000 > 48
#include <opm/simulators/linalg/bda/BdaBridge.hpp>
#include <opm/simulators/linalg/bda/WellContributions.hpp>
#include <dune/common/fvector.hh>
#include <dune/istl/bvector.hh>
+1
View File
@@ -30,6 +30,7 @@
BOOST_VERSION / 100 % 1000 > 48
#include <opm/simulators/linalg/bda/BdaBridge.hpp>
#include <opm/simulators/linalg/bda/WellContributions.hpp>
#include <dune/common/fvector.hh>
#include <dune/istl/bvector.hh>