diff --git a/opm/simulators/linalg/bda/rocm/hipKernels.cpp b/opm/simulators/linalg/bda/rocm/hipKernels.cpp index 6f5683059..892d1e090 100644 --- a/opm/simulators/linalg/bda/rocm/hipKernels.cpp +++ b/opm/simulators/linalg/bda/rocm/hipKernels.cpp @@ -304,12 +304,12 @@ init(int verbosity_) template void HipKernels:: -vmul(const Scalar alpha, - Scalar* in1, - Scalar* in2, - Scalar* out, - int N, - hipStream_t stream) +vmul([[maybe_unused]] const Scalar alpha, + [[maybe_unused]] Scalar* in1, + [[maybe_unused]] Scalar* in2, + [[maybe_unused]] Scalar* out, + [[maybe_unused]] int N, + [[maybe_unused]] hipStream_t stream) { Timer t_vmul; #ifdef __HIP__ @@ -333,11 +333,11 @@ vmul(const Scalar alpha, template void HipKernels:: -full_to_pressure_restriction(const Scalar* fine_y, - Scalar* weights, - Scalar* coarse_y, - int Nb, - hipStream_t stream) +full_to_pressure_restriction([[maybe_unused]] const Scalar* fine_y, + [[maybe_unused]] Scalar* weights, + [[maybe_unused]] Scalar* coarse_y, + [[maybe_unused]] int Nb, + [[maybe_unused]] hipStream_t stream) { Timer t; #ifdef __HIP__ @@ -361,11 +361,11 @@ full_to_pressure_restriction(const Scalar* fine_y, template void HipKernels:: -add_coarse_pressure_correction(Scalar* coarse_x, - Scalar* fine_x, - int pressure_idx, - int Nb, - hipStream_t stream) +add_coarse_pressure_correction([[maybe_unused]] Scalar* coarse_x, + [[maybe_unused]] Scalar* fine_x, + [[maybe_unused]] int pressure_idx, + [[maybe_unused]] int Nb, + [[maybe_unused]] hipStream_t stream) { Timer t; #ifdef __HIP__ @@ -389,11 +389,11 @@ add_coarse_pressure_correction(Scalar* coarse_x, template void HipKernels:: -prolongate_vector(const Scalar* in, - Scalar* out, - const int* cols, - int N, - hipStream_t stream) +prolongate_vector([[maybe_unused]] const Scalar* in, + [[maybe_unused]] Scalar* out, + [[maybe_unused]] const int* cols, + [[maybe_unused]] int N, + [[maybe_unused]] hipStream_t stream) { Timer t; @@ -419,15 +419,15 @@ prolongate_vector(const Scalar* in, template void HipKernels:: -residual(Scalar* vals, - int* cols, - int* rows, - Scalar* x, - const Scalar* rhs, - Scalar* out, - int Nb, - unsigned int block_size, - hipStream_t stream) +residual([[maybe_unused]] Scalar* vals, + [[maybe_unused]] int* cols, + [[maybe_unused]] int* rows, + [[maybe_unused]] Scalar* x, + [[maybe_unused]] const Scalar* rhs, + [[maybe_unused]] Scalar* out, + [[maybe_unused]] int Nb, + [[maybe_unused]] unsigned int block_size, + [[maybe_unused]] hipStream_t stream) { Timer t_residual; @@ -458,14 +458,14 @@ residual(Scalar* vals, template void HipKernels:: -spmv(Scalar* vals, - int* cols, - int* rows, - Scalar* x, - Scalar* y, - int Nb, - unsigned int block_size, - hipStream_t stream) +spmv([[maybe_unused]] Scalar* vals, + [[maybe_unused]] int* cols, + [[maybe_unused]] int* rows, + [[maybe_unused]] Scalar* x, + [[maybe_unused]] Scalar* y, + [[maybe_unused]] int Nb, + [[maybe_unused]] unsigned int block_size, + [[maybe_unused]] hipStream_t stream) {//NOTE: block_size not used since I use this kernel only for block sizes 1, other uses use rocsparse! Timer t_spmv; #ifdef __HIP__ diff --git a/opm/simulators/linalg/bda/rocm/rocsparseBILU0.cpp b/opm/simulators/linalg/bda/rocm/rocsparseBILU0.cpp index 9f48ed77b..bd4a721f5 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparseBILU0.cpp +++ b/opm/simulators/linalg/bda/rocm/rocsparseBILU0.cpp @@ -92,8 +92,8 @@ analyze_matrix(BlockedMatrix *mat) { template bool rocsparseBILU0:: -analyze_matrix(BlockedMatrix *mat, - BlockedMatrix *jacMat) +analyze_matrix(BlockedMatrix*, + BlockedMatrix*) { std::size_t d_bufferSize_M, d_bufferSize_L, d_bufferSize_U, d_bufferSize; Timer t; @@ -163,8 +163,8 @@ create_preconditioner(BlockedMatrix *mat) { template bool rocsparseBILU0:: -create_preconditioner(BlockedMatrix *mat, - BlockedMatrix *jacMat) +create_preconditioner(BlockedMatrix*, + BlockedMatrix*) { Timer t; bool result = true; @@ -253,7 +253,6 @@ update_system_on_gpu(Scalar *d_Avals) { template void rocsparseBILU0:: apply(Scalar& y, Scalar& x) { - Scalar zero = 0.0; Scalar one = 1.0; Timer t_apply; @@ -274,7 +273,7 @@ apply(Scalar& y, Scalar& x) { } } -#define INSTANCE_TYPE(T) \ +#define INSTANTIATE_TYPE(T) \ template class rocsparseBILU0; \ template class rocsparseBILU0; \ template class rocsparseBILU0; \ @@ -282,5 +281,6 @@ apply(Scalar& y, Scalar& x) { template class rocsparseBILU0; \ template class rocsparseBILU0; -INSTANCE_TYPE(double) +INSTANTIATE_TYPE(double) + } // namespace Opm diff --git a/opm/simulators/linalg/bda/rocm/rocsparseBILU0.hpp b/opm/simulators/linalg/bda/rocm/rocsparseBILU0.hpp index 496d9dff6..011d51b5e 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparseBILU0.hpp +++ b/opm/simulators/linalg/bda/rocm/rocsparseBILU0.hpp @@ -20,9 +20,6 @@ #ifndef OPM_ROCSPARSEBILU0_HPP #define OPM_ROCSPARSEBILU0_HPP -#include -#include - #include #include @@ -70,17 +67,17 @@ public: bool initialize(std::shared_ptr> matrix, std::shared_ptr> jacMatrix, rocsparse_int *d_Arows, - rocsparse_int *d_Acols); + rocsparse_int *d_Acols) override; /// Analysis, extract parallelism if specified /// \param[in] mat matrix A - bool analyze_matrix(BlockedMatrix *mat); + bool analyze_matrix(BlockedMatrix *mat) override; /// Analysis, extract parallelism if specified /// \param[in] mat matrix A /// \param[in] jacMat matrix for preconditioner, analyze this as well bool analyze_matrix(BlockedMatrix *mat, - BlockedMatrix *jacMat); + BlockedMatrix *jacMat) override; /// ILU decomposition /// \param[in] mat matrix A to decompose @@ -102,13 +99,12 @@ public: #if HAVE_OPENCL // apply preconditioner, x = prec(y) - void apply(const cl::Buffer& y, - cl::Buffer& x) {} + void apply(const cl::Buffer&, cl::Buffer&) override {} #endif /// Copy matrix A values to GPU /// \param[in] mVals Input values - void copy_system_to_gpu(Scalar *mVals); + void copy_system_to_gpu(Scalar *mVals) override; /// Reassign pointers, in case the addresses of the Dune variables have changed --> TODO: check when/if we need this method // /// \param[in] vals New values @@ -117,7 +113,7 @@ public: /// Update GPU values after a new assembly is done /// \param[in] b New b vector - void update_system_on_gpu(Scalar *b); + void update_system_on_gpu(Scalar *b) override; }; } // namespace Opm diff --git a/opm/simulators/linalg/bda/rocm/rocsparseCPR.cpp b/opm/simulators/linalg/bda/rocm/rocsparseCPR.cpp index 5a95dd475..343a465a0 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparseCPR.cpp +++ b/opm/simulators/linalg/bda/rocm/rocsparseCPR.cpp @@ -95,14 +95,14 @@ analyze_matrix(BlockedMatrix *mat_) { template bool rocsparseCPR:: analyze_matrix(BlockedMatrix *mat_, - BlockedMatrix *jacMat) + BlockedMatrix *jacMat_) { this->Nb = mat_->Nb; this->nnzb = mat_->nnzbs; this->N = Nb * block_size; this->nnz = nnzb * block_size * block_size; - bool success = bilu0->analyze_matrix(mat_, jacMat); + bool success = bilu0->analyze_matrix(mat_, jacMat_); this->mat = mat_; return success; @@ -111,10 +111,10 @@ analyze_matrix(BlockedMatrix *mat_, template bool rocsparseCPR:: create_preconditioner(BlockedMatrix *mat_, - BlockedMatrix *jacMat) + BlockedMatrix *jacMat_) { Dune::Timer t_bilu0; - bool result = bilu0->create_preconditioner(mat_, jacMat); + bool result = bilu0->create_preconditioner(mat_, jacMat_); if (verbosity >= 3) { std::ostringstream out; out << "rocsparseCPR create_preconditioner bilu0(): " << t_bilu0.stop() << " s"; @@ -223,8 +223,6 @@ amg_cycle_gpu(const int level, RocmMatrix *A = &d_Amatrices[level]; RocmMatrix *R = &d_Rmatrices[level]; int Ncur = A->Nb; - Scalar zero = 0.0; - Scalar one = 1.0; rocsparse_mat_info spmv_info; rocsparse_mat_descr descr_R; @@ -324,7 +322,7 @@ apply(Scalar& y, } } -#define INSTANCE_TYPE(T) \ +#define INSTANTIATE_TYPE(T) \ template class rocsparseCPR; \ template class rocsparseCPR; \ template class rocsparseCPR; \ @@ -332,7 +330,6 @@ apply(Scalar& y, template class rocsparseCPR; \ template class rocsparseCPR; -INSTANCE_TYPE(double) +INSTANTIATE_TYPE(double) + } // namespace Opm - - diff --git a/opm/simulators/linalg/bda/rocm/rocsparseCPR.hpp b/opm/simulators/linalg/bda/rocm/rocsparseCPR.hpp index 72459c1f7..ffbea9023 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparseCPR.hpp +++ b/opm/simulators/linalg/bda/rocm/rocsparseCPR.hpp @@ -85,28 +85,28 @@ public: bool initialize(std::shared_ptr> matrix, std::shared_ptr> jacMatrix, rocsparse_int *d_Arows, - rocsparse_int *d_Acols); + rocsparse_int *d_Acols) override; /// Analysis, extract parallelism if specified /// \param[in] mat matrix A - bool analyze_matrix(BlockedMatrix *mat); + bool analyze_matrix(BlockedMatrix *mat) override; /// Analysis, extract parallelism if specified /// \param[in] mat matrix A /// \param[in] jacMat matrix for preconditioner, analyze this as well bool analyze_matrix(BlockedMatrix *mat, - BlockedMatrix *jacMat); + BlockedMatrix *jacMat) override; /// Create AMG preconditioner and perform ILU decomposition /// \param[in] mat matrix A - bool create_preconditioner(BlockedMatrix *mat); + bool create_preconditioner(BlockedMatrix *mat) override; /// Create AMG preconditioner and perform ILU decomposition /// \param[in] mat matrix A /// \param[in] jacMat matrix for preconditioner, decompose this one if used bool create_preconditioner(BlockedMatrix *mat, - BlockedMatrix *jacMat); + BlockedMatrix *jacMat) override; /// Apply preconditioner, x = prec(y) /// applies blocked ilu0 @@ -118,13 +118,12 @@ public: #if HAVE_OPENCL // apply preconditioner, x = prec(y) - void apply(const cl::Buffer& y, - cl::Buffer& x) {} + void apply(const cl::Buffer&, cl::Buffer&) override {} #endif /// Copy matrix A values to GPU /// \param[in] mVals Input values - void copy_system_to_gpu(Scalar *b); + void copy_system_to_gpu(Scalar *b) override; /// Reassign pointers, in case the addresses of the Dune variables have changed --> TODO: check when/if we need this method /// \param[in] vals array of nonzeroes, each block is stored row-wise and contiguous, contains nnz values @@ -133,7 +132,7 @@ public: /// Update linear system to GPU /// \param[in] b input vector, contains N values - void update_system_on_gpu(Scalar *b); + void update_system_on_gpu(Scalar *b) override; }; } // namespace Opm diff --git a/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.cpp b/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.cpp index f5449b903..17e44e9a7 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.cpp +++ b/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.cpp @@ -38,38 +38,39 @@ create(PreconditionerType type, return std::make_unique >(verbosity); case PreconditionerType::CPR: return std::make_unique >(verbosity); + default: + OPM_THROW(std::logic_error, + "Invalid preconditioner type " + std::to_string(static_cast(type))); } - - OPM_THROW(std::logic_error, - "Invalid preconditioner type " + std::to_string(static_cast(type))); } template void rocsparsePreconditioner:: -set_matrix_analysis(rocsparse_mat_descr descr_L, - rocsparse_mat_descr descr_U) +set_matrix_analysis(rocsparse_mat_descr desc_L, + rocsparse_mat_descr desc_U) { - descr_L = descr_L; - descr_U = descr_U; + descr_L = desc_L; + descr_U = desc_U; } template void rocsparsePreconditioner:: -set_context(rocsparse_handle handle, - rocsparse_direction dir, - rocsparse_operation operation, - hipStream_t stream) +set_context(rocsparse_handle handle_, + rocsparse_direction dir_, + rocsparse_operation operation_, + hipStream_t stream_) { - this->handle = handle; - this->dir = dir; - this->operation = operation; - this->stream = stream; + this->handle = handle_; + this->dir = dir_; + this->operation = operation_; + this->stream = stream_; } template void rocsparsePreconditioner:: -setJacMat(BlockedMatrix jacMat) { - this->jacMat = std::make_shared>(jacMat); +setJacMat(const BlockedMatrix& jMat) +{ + this->jacMat = std::make_shared>(jMat); } #define INSTANTIATE_TYPE(T) \ diff --git a/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.hpp b/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.hpp index ef50b3eeb..7a4ad28a7 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.hpp +++ b/opm/simulators/linalg/bda/rocm/rocsparsePreconditioner.hpp @@ -84,7 +84,7 @@ public: rocsparse_operation operation, hipStream_t stream); - void setJacMat(BlockedMatrix jacMat); + void setJacMat(const BlockedMatrix& jacMat); }; } //namespace Opm diff --git a/opm/simulators/linalg/bda/rocm/rocsparseSolverBackend.cpp b/opm/simulators/linalg/bda/rocm/rocsparseSolverBackend.cpp index 3a1781076..c729cb27f 100644 --- a/opm/simulators/linalg/bda/rocm/rocsparseSolverBackend.cpp +++ b/opm/simulators/linalg/bda/rocm/rocsparseSolverBackend.cpp @@ -63,18 +63,16 @@ rocsparseSolverBackend(int verbosity_, int maxit_, Scalar tolerance_, : Base(verbosity_, maxit_, tolerance_, platformID_, deviceID_) { int numDevices = 0; - bool use_cpr, use_isai; + bool use_cpr; if (linsolver.compare("ilu0") == 0) { use_cpr = false; - use_isai = false; } else if (linsolver.compare("cpr_quasiimpes") == 0) { use_cpr = true; - use_isai = false; } else if (linsolver.compare("isai") == 0) { - OPM_THROW(std::logic_error, "Error rocsparseSolver does not support --linerar-solver=isai"); + OPM_THROW(std::logic_error, "Error rocsparseSolver does not support --linear-solver=isai"); } else if (linsolver.compare("cpr_trueimpes") == 0) { - OPM_THROW(std::logic_error, "Error rocsparseSolver does not support --linerar-solver=cpr_trueimpes"); + OPM_THROW(std::logic_error, "Error rocsparseSolver does not support --linear-solver=cpr_trueimpes"); } else { OPM_THROW(std::logic_error, "Error unknown value for argument --linear-solver, " + linsolver); } @@ -100,11 +98,11 @@ rocsparseSolverBackend(int verbosity_, int maxit_, Scalar tolerance_, ROCSPARSE_CHECK(rocsparse_set_stream(handle, stream)); ROCBLAS_CHECK(rocblas_set_stream(blas_handle, stream)); - using PreconditionerType = typename Opm::Accelerator::PreconditionerType; + using PCType = typename Opm::Accelerator::PreconditionerType; if (use_cpr) { - prec = rocsparsePreconditioner::create(PreconditionerType::CPR, verbosity); + prec = rocsparsePreconditioner::create(PCType::CPR, verbosity); } else { - prec = rocsparsePreconditioner::create(PreconditionerType::BILU0, verbosity); + prec = rocsparsePreconditioner::create(PCType::BILU0, verbosity); } prec->set_context(handle, dir, operation, stream);