From b2ca7e8769e1735b21e9b5162d7568e26ef28363 Mon Sep 17 00:00:00 2001 From: Kjetil Olsen Lye Date: Thu, 30 Apr 2026 20:50:31 +0000 Subject: [PATCH] GPU portability: OPM_HOST_DEVICE/OPM_THROW additions and fluidSystem() use Decorates inline helpers in blackoil*modules.hh, blackoilprimaryvariables.hh, directionalmobility.hh, tpfalinearizer.hh, NewTranFluxModule.hpp, TemperatureModel.hpp and BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp with OPM_HOST_DEVICE, replaces a few raw `throw`s with OPM_THROW, and switches static `FluidSystem::` accesses to instance-based `fluidState.fluidSystem()` / `getFluidSystem()` so the same code works when the fluid system lives on the device. Test test_gpu_linear_two_phase_material.cu picks up the matching template adjustments. --- ...ilEnergyIntensiveQuantitiesGlobalIndex.hpp | 86 +++---- opm/models/blackoil/blackoilbrinemodules.hh | 60 ++--- opm/models/blackoil/blackoilenergymodules.hh | 214 +++++++++--------- opm/models/blackoil/blackoilfoammodules.hh | 3 +- .../blackoil/blackoilprimaryvariables.hh | 10 +- opm/models/common/directionalmobility.hh | 11 +- .../discretization/common/tpfalinearizer.hh | 4 +- opm/simulators/flow/NewTranFluxModule.hpp | 7 +- opm/simulators/flow/TemperatureModel.hpp | 21 +- .../test_gpu_linear_two_phase_material.cu | 10 +- 10 files changed, 222 insertions(+), 204 deletions(-) diff --git a/flowexperimental/BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp b/flowexperimental/BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp index bcb684800d..0c62cbbb26 100644 --- a/flowexperimental/BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp +++ b/flowexperimental/BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp @@ -59,20 +59,20 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex; public: - void updateTemperature_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalSpaceIdx, - [[maybe_unused]] unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalSpaceIdx, + [[maybe_unused]] unsigned timeIdx) { auto& fs = this->asImp_().fluidState_; Scalar T = problem.temperature(globalSpaceIdx, timeIdx); fs.setTemperature(T); } - void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalSpaceIdx, - [[maybe_unused]] unsigned timeIdx, - const typename FluidSystem::template ParameterCache&) + OPM_HOST_DEVICE void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalSpaceIdx, + [[maybe_unused]] unsigned timeIdx, + const typename FluidSystem::template ParameterCache&) { } }; @@ -185,17 +187,17 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex; public: - void updateTemperature_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalSpaceIdx, - [[maybe_unused]] unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalSpaceIdx, + [[maybe_unused]] unsigned timeIdx) { } - void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalSpaceIdx, - [[maybe_unused]] unsigned timeIdx, - const typename FluidSystem::template ParameterCache&) + OPM_HOST_DEVICE void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalSpaceIdx, + [[maybe_unused]] unsigned timeIdx, + const typename FluidSystem::template ParameterCache&) { } }; diff --git a/opm/models/blackoil/blackoilbrinemodules.hh b/opm/models/blackoil/blackoilbrinemodules.hh index caa5daf74c..6947429d8a 100644 --- a/opm/models/blackoil/blackoilbrinemodules.hh +++ b/opm/models/blackoil/blackoilbrinemodules.hh @@ -98,7 +98,7 @@ public: static void registerParameters() {} - static bool primaryVarApplies(unsigned pvIdx) + static OPM_HOST_DEVICE bool primaryVarApplies(unsigned pvIdx) { if constexpr (enableBrine) { return pvIdx == saltConcentrationIdx; @@ -135,7 +135,7 @@ public: return static_cast(1.0); } - static bool eqApplies(unsigned eqIdx) + static OPM_HOST_DEVICE bool eqApplies(unsigned eqIdx) { if constexpr (enableBrine) { return eqIdx == contiBrineEqIdx; @@ -405,18 +405,18 @@ public: * primary variables * */ - void updateSaltConcentration_(const ElementContext& elemCtx, - unsigned dofIdx, - unsigned timeIdx) + OPM_HOST_DEVICE void updateSaltConcentration_(const ElementContext& elemCtx, + unsigned dofIdx, + unsigned timeIdx) { const PrimaryVariables& priVars = elemCtx.primaryVars(dofIdx, timeIdx); const LinearizationType lintype = elemCtx.linearizationType(); updateSaltConcentration_(priVars, timeIdx, lintype); } - void updateSaltConcentration_(const PrimaryVariables& priVars, - const unsigned timeIdx, - const LinearizationType lintype) + OPM_HOST_DEVICE void updateSaltConcentration_(const PrimaryVariables& priVars, + const unsigned timeIdx, + const LinearizationType lintype) { const unsigned pvtnumRegionIdx = priVars.pvtRegionIndex(); auto& fs = asImp_().fluidState_; @@ -442,9 +442,9 @@ public: } } - void saltPropertiesUpdate_([[maybe_unused]] const ElementContext& elemCtx, - [[maybe_unused]] unsigned dofIdx, - [[maybe_unused]] unsigned timeIdx) + OPM_HOST_DEVICE void saltPropertiesUpdate_([[maybe_unused]] const ElementContext& elemCtx, + [[maybe_unused]] unsigned dofIdx, + [[maybe_unused]] unsigned timeIdx) { if constexpr (enableSaltPrecipitation) { const Evaluation porosityFactor = min(1.0 - asImp_().fluidState_.saltSaturation(), 1.0); //phi/phi_0 @@ -455,20 +455,20 @@ public: } } - const Evaluation& brineRefDensity() const + OPM_HOST_DEVICE const Evaluation& brineRefDensity() const { return refDensity_; } - Scalar saltSolubility() const + OPM_HOST_DEVICE Scalar saltSolubility() const { return saltSolubility_; } - Scalar saltDensity() const + OPM_HOST_DEVICE Scalar saltDensity() const { return saltDensity_; } - const Evaluation& permFactor() const + OPM_HOST_DEVICE const Evaluation& permFactor() const { return permFactor_; } protected: - Implementation& asImp_() + OPM_HOST_DEVICE Implementation& asImp_() { return *static_cast(this); } Evaluation saltConcentration_; @@ -487,27 +487,27 @@ class BlackOilBrineIntensiveQuantities using Scalar = GetPropType; public: - void updateSaltConcentration_(const ElementContext&, - unsigned, - unsigned) + OPM_HOST_DEVICE void updateSaltConcentration_(const ElementContext&, + unsigned, + unsigned) {} - void saltPropertiesUpdate_(const ElementContext&, - unsigned, - unsigned) + OPM_HOST_DEVICE void saltPropertiesUpdate_(const ElementContext&, + unsigned, + unsigned) {} - const Evaluation& brineRefDensity() const - { throw std::runtime_error("brineRefDensity() called but brine are disabled"); } + OPM_HOST_DEVICE const Evaluation& brineRefDensity() const + { OPM_THROW(std::runtime_error, "brineRefDensity() called but brine are disabled"); } - const Scalar saltSolubility() const - { throw std::logic_error("saltSolubility() called but salt precipitation is disabled"); } + OPM_HOST_DEVICE const Scalar saltSolubility() const + { OPM_THROW(std::logic_error, "saltSolubility() called but salt precipitation is disabled"); } - const Scalar saltDensity() const - { throw std::logic_error("saltDensity() called but salt precipitation is disabled"); } + OPM_HOST_DEVICE const Scalar saltDensity() const + { OPM_THROW(std::logic_error, "saltDensity() called but salt precipitation is disabled"); } - const Evaluation& permFactor() const - { throw std::logic_error("permFactor() called but salt precipitation is disabled"); } + OPM_HOST_DEVICE const Evaluation& permFactor() const + { OPM_THROW(std::logic_error, "permFactor() called but salt precipitation is disabled"); } }; } // namespace Opm diff --git a/opm/models/blackoil/blackoilenergymodules.hh b/opm/models/blackoil/blackoilenergymodules.hh index 5d1aebf2fa..e2cce9cd2a 100644 --- a/opm/models/blackoil/blackoilenergymodules.hh +++ b/opm/models/blackoil/blackoilenergymodules.hh @@ -30,6 +30,7 @@ #include +#include #include #include @@ -104,7 +105,7 @@ public: } } - static bool primaryVarApplies(unsigned pvIdx) + static OPM_HOST_DEVICE bool primaryVarApplies(unsigned pvIdx) { if constexpr (enableFullyImplicitThermal) { return pvIdx == temperatureIdx; @@ -129,7 +130,7 @@ public: return static_cast(1.0); } - static bool eqApplies(unsigned eqIdx) + static OPM_HOST_DEVICE bool eqApplies(unsigned eqIdx) { if constexpr (enableFullyImplicitThermal) { return eqIdx == contiEnergyEqIdx; @@ -241,11 +242,11 @@ public: } template - static void addPhaseEnthalpyFlux_(RateVector& flux, - unsigned phaseIdx, - const ElementContext& elemCtx, - unsigned scvfIdx, - unsigned timeIdx) + OPM_HOST_DEVICE static void addPhaseEnthalpyFlux_(RateVector& flux, + unsigned phaseIdx, + const ElementContext& elemCtx, + unsigned scvfIdx, + unsigned timeIdx) { const auto& extQuants = elemCtx.extensiveQuantities(scvfIdx, timeIdx); const unsigned upIdx = extQuants.upstreamIndex(phaseIdx); @@ -258,8 +259,8 @@ public: fs); } - static void addToEnthalpyRate(RateVector& flux, - const Evaluation& hRate) + OPM_HOST_DEVICE static void addToEnthalpyRate(RateVector& flux, + const Evaluation& hRate) { if constexpr (enableFullyImplicitThermal) { flux[contiEnergyEqIdx] += hRate; @@ -270,8 +271,8 @@ public: * \brief Assign the energy specific primary variables to a PrimaryVariables object */ template - static void assignPrimaryVars(PrimaryVariables& priVars, - const FluidState& fluidState) + OPM_HOST_DEVICE static void assignPrimaryVars(PrimaryVariables& priVars, + const FluidState& fluidState) { if constexpr (enableFullyImplicitThermal) { priVars[temperatureIdx] = getValue(fluidState.temperature(/*phaseIdx=*/0)); @@ -281,9 +282,9 @@ public: /*! * \brief Do a Newton-Raphson update the primary variables of the energys. */ - static void updatePrimaryVars(PrimaryVariables& newPv, - const PrimaryVariables& oldPv, - const EqVector& delta) + OPM_HOST_DEVICE static void updatePrimaryVars(PrimaryVariables& newPv, + const PrimaryVariables& oldPv, + const EqVector& delta) { if constexpr (enableFullyImplicitThermal) { // do a plain unchopped Newton update @@ -294,8 +295,8 @@ public: /*! * \brief Return how much a Newton-Raphson update is considered an error */ - static Scalar computeUpdateError(const PrimaryVariables&, - const EqVector&) + OPM_HOST_DEVICE static Scalar computeUpdateError(const PrimaryVariables&, + const EqVector&) { // do not consider consider the cange of energy primary variables for // convergence @@ -306,7 +307,7 @@ public: /*! * \brief Return how much a residual is considered an error */ - static Scalar computeResidualError(const EqVector& resid) + OPM_HOST_DEVICE static Scalar computeResidualError(const EqVector& resid) { // do not weight the residual of energy when it comes to convergence return std::abs(scalarValue(resid[contiEnergyEqIdx])); @@ -385,9 +386,9 @@ public: * \brief Update the temperature of the intensive quantity's fluid state * */ - void updateTemperature_(const ElementContext& elemCtx, - unsigned dofIdx, - unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_(const ElementContext& elemCtx, + unsigned dofIdx, + unsigned timeIdx) { auto& fs = asImp_().fluidState_; const auto& priVars = elemCtx.primaryVars(dofIdx, timeIdx); @@ -400,11 +401,11 @@ public: * \brief Update the temperature of the intensive quantity's fluid state * */ - void updateTemperature_([[maybe_unused]] const Problem& problem, - const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalDofIdx, - const unsigned timeIdx, - const LinearizationType& lintype) + OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem, + const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalDofIdx, + const unsigned timeIdx, + const LinearizationType& lintype) { auto& fs = asImp_().fluidState_; fs.setTemperature(priVars.makeEvaluation(temperatureIdx, timeIdx, lintype)); @@ -414,27 +415,27 @@ public: * \brief Compute the intensive quantities needed to handle energy conservation * */ - void updateEnergyQuantities_(const ElementContext& elemCtx, - unsigned dofIdx, - unsigned timeIdx) + OPM_HOST_DEVICE void updateEnergyQuantities_(const ElementContext& elemCtx, + unsigned dofIdx, + unsigned timeIdx) { updateEnergyQuantities_(elemCtx.problem(), elemCtx.globalSpaceIndex(dofIdx, timeIdx), timeIdx); } - void updateEnergyQuantities_(const Problem& problem, - const unsigned globalSpaceIdx, - const unsigned timeIdx) + OPM_HOST_DEVICE void updateEnergyQuantities_(const Problem& problem, + const unsigned globalSpaceIdx, + const unsigned timeIdx) { auto& fs = asImp_().fluidState_; // compute the specific enthalpy of the fluids, the specific enthalpy of the rock // and the thermal conductivity coefficients for (int phaseIdx = 0; phaseIdx < numPhases; ++ phaseIdx) { - if (!FluidSystem::phaseIsActive(phaseIdx)) { + if (!asImp_().getFluidSystem().phaseIsActive(phaseIdx)) { continue; } - const auto& h = FluidSystem::enthalpy(fs, phaseIdx, problem.pvtRegionIndex(globalSpaceIdx)); + const auto& h = asImp_().getFluidSystem().enthalpy(fs, phaseIdx, problem.pvtRegionIndex(globalSpaceIdx)); fs.setEnthalpy(phaseIdx, h); } @@ -462,7 +463,7 @@ public: { return rockFraction_; } protected: - Implementation& asImp_() + OPM_HOST_DEVICE Implementation& asImp_() { return *static_cast(this); } Evaluation rockInternalEnergy_; @@ -484,51 +485,54 @@ class BlackOilEnergyIntensiveQuantities - void updateTemperature_(const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - unsigned globalDofIdx, - unsigned timeIdx, - [[maybe_unused]] const LinearizationType& lintype - ) + OPM_HOST_DEVICE void updateTemperature_(const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + unsigned globalDofIdx, + unsigned timeIdx, + [[maybe_unused]] const LinearizationType& lintype) { updateTemperature_(problem, globalDofIdx, timeIdx); } - void updateTemperature_(const Problem& problem, unsigned globalDofIdx, unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_(const Problem& problem, + unsigned globalDofIdx, + unsigned timeIdx) { auto& fs = asImp_().fluidState_; const Scalar T = problem.temperature(globalDofIdx, timeIdx); fs.setTemperature(T); } - void updateEnergyQuantities_(const ElementContext&, - unsigned, - unsigned, - const typename FluidSystem::template ParameterCache&) + OPM_HOST_DEVICE void updateEnergyQuantities_(const ElementContext&, + unsigned, + unsigned, + const typename FluidSystem::template ParameterCache&) {} - const Evaluation& rockInternalEnergy() const + OPM_HOST_DEVICE const Evaluation& rockInternalEnergy() const { - throw std::logic_error("Requested the rock internal energy, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the rock internal energy, which is " + "unavailable because energy is not conserved"); } - const Evaluation& totalThermalConductivity() const + OPM_HOST_DEVICE const Evaluation& totalThermalConductivity() const { - throw std::logic_error("Requested the total thermal conductivity, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the total thermal conductivity, which is " + "unavailable because energy is not conserved"); } protected: - Implementation& asImp_() + OPM_HOST_DEVICE Implementation& asImp_() { return *static_cast(this); } }; @@ -548,26 +552,28 @@ class BlackOilEnergyIntensiveQuantities - void updateTemperature_(const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - unsigned globalDofIdx, - unsigned timeIdx, - [[maybe_unused]] const LinearizationType& lintype) + OPM_HOST_DEVICE void updateTemperature_(const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + unsigned globalDofIdx, + unsigned timeIdx, + [[maybe_unused]] const LinearizationType& lintype) { updateTemperature_(problem, globalDofIdx, timeIdx); } @@ -576,32 +582,34 @@ public: * \brief Compute the intensive quantities needed to handle energy conservation * */ - void updateEnergyQuantities_([[maybe_unused]] const ElementContext& elemCtx, - [[maybe_unused]] unsigned dofIdx, - [[maybe_unused]] unsigned timeIdx) + OPM_HOST_DEVICE void updateEnergyQuantities_([[maybe_unused]] const ElementContext& elemCtx, + [[maybe_unused]] unsigned dofIdx, + [[maybe_unused]] unsigned timeIdx) { } - void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const unsigned globalSpaceIdx, - [[maybe_unused]] const unsigned timeIdx) + OPM_HOST_DEVICE void updateEnergyQuantities_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const unsigned globalSpaceIdx, + [[maybe_unused]] const unsigned timeIdx) { } - const Evaluation& rockInternalEnergy() const + OPM_HOST_DEVICE const Evaluation& rockInternalEnergy() const { - throw std::logic_error("Requested the rock internal energy, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the rock internal energy, which is " + "unavailable because energy is not conserved"); } - const Evaluation& totalThermalConductivity() const + OPM_HOST_DEVICE const Evaluation& totalThermalConductivity() const { - throw std::logic_error("Requested the total thermal conductivity, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the total thermal conductivity, which is " + "unavailable because energy is not conserved"); } protected: - Implementation& asImp_() + OPM_HOST_DEVICE Implementation& asImp_() { return *static_cast(this); } }; @@ -617,41 +625,43 @@ class BlackOilEnergyIntensiveQuantities using Problem = GetPropType; public: - void updateTemperature_([[maybe_unused]] const ElementContext& elemCtx, - [[maybe_unused]] unsigned dofIdx, - [[maybe_unused]] unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const ElementContext& elemCtx, + [[maybe_unused]] unsigned dofIdx, + [[maybe_unused]] unsigned timeIdx) { } template - void updateTemperature_([[maybe_unused]] const Problem& problem, - [[maybe_unused]] const PrimaryVariables& priVars, - [[maybe_unused]] unsigned globalDofIdx, - [[maybe_unused]] unsigned timeIdx, - [[maybe_unused]] const LinearizationType& lintype) + OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem, + [[maybe_unused]] const PrimaryVariables& priVars, + [[maybe_unused]] unsigned globalDofIdx, + [[maybe_unused]] unsigned timeIdx, + [[maybe_unused]] const LinearizationType& lintype) { } - void updateEnergyQuantities_(const ElementContext&, - unsigned, - unsigned, - const typename FluidSystem::template ParameterCache&) + OPM_HOST_DEVICE void updateEnergyQuantities_(const ElementContext&, + unsigned, + unsigned, + const typename FluidSystem::template ParameterCache&) {} - const Evaluation& rockInternalEnergy() const + OPM_HOST_DEVICE const Evaluation& rockInternalEnergy() const { - throw std::logic_error("Requested the rock internal energy, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the rock internal energy, which is " + "unavailable because energy is not conserved"); } - const Evaluation& totalThermalConductivity() const + OPM_HOST_DEVICE const Evaluation& totalThermalConductivity() const { - throw std::logic_error("Requested the total thermal conductivity, which is " - "unavailable because energy is not conserved"); + OPM_THROW(std::logic_error, + "Requested the total thermal conductivity, which is " + "unavailable because energy is not conserved"); } protected: - Implementation& asImp_() + OPM_HOST_DEVICE Implementation& asImp_() { return *static_cast(this); } }; @@ -873,8 +883,8 @@ public: const BoundaryFluidState& /*boundaryFs*/) {} - const Evaluation& energyFlux() const - { throw std::logic_error("Requested the energy flux, but energy is not conserved"); } + OPM_HOST_DEVICE const Evaluation& energyFlux() const + { OPM_THROW(std::logic_error, "Requested the energy flux, but energy is not conserved"); } }; template @@ -934,8 +944,8 @@ public: const BoundaryFluidState& /*boundaryFs*/) { } - const Evaluation& energyFlux() const - { throw std::logic_error("Requested the energy flux, but energy is not conserved"); } + OPM_HOST_DEVICE const Evaluation& energyFlux() const + { OPM_THROW(std::logic_error, "Requested the energy flux, but energy is not conserved"); } }; template @@ -982,8 +992,8 @@ public: const BoundaryFluidState& /*boundaryFs*/) {} - const Evaluation& energyFlux() const - { throw std::logic_error("Requested the energy flux, but energy is not conserved"); } + OPM_HOST_DEVICE const Evaluation& energyFlux() const + { OPM_THROW(std::logic_error, "Requested the energy flux, but energy is not conserved"); } }; } // namespace Opm diff --git a/opm/models/blackoil/blackoilfoammodules.hh b/opm/models/blackoil/blackoilfoammodules.hh index d90f453aff..7fb156b091 100644 --- a/opm/models/blackoil/blackoilfoammodules.hh +++ b/opm/models/blackoil/blackoilfoammodules.hh @@ -30,6 +30,7 @@ #include +#include #include #include @@ -184,7 +185,7 @@ public: Toolbox::template decay(intQuants.solventInverseFormationVolumeFactor()); } } else { - throw std::runtime_error("Transport phase is GAS/WATER/SOLVENT"); + OPM_THROW(std::runtime_error, "Transport phase is GAS/WATER/SOLVENT"); } // Avoid singular matrix if no gas is present. diff --git a/opm/models/blackoil/blackoilprimaryvariables.hh b/opm/models/blackoil/blackoilprimaryvariables.hh index 1ed590b8c3..8eea039647 100644 --- a/opm/models/blackoil/blackoilprimaryvariables.hh +++ b/opm/models/blackoil/blackoilprimaryvariables.hh @@ -1024,10 +1024,10 @@ private: template OPM_HOST_DEVICE void computeCapillaryPressures_(Container& result, - Scalar so, - Scalar sg, - Scalar sw, - const MaterialLawParams& matParams) const + Scalar so, + Scalar sg, + Scalar sw, + const MaterialLawParams& matParams) const { using SatOnlyFluidState = SimpleModularFluidStategetPressureScale()); } // NOTE: When adding new member variables, be sure to update the template copy constructor, diff --git a/opm/models/common/directionalmobility.hh b/opm/models/common/directionalmobility.hh index bcf2056cbe..35f6a81865 100644 --- a/opm/models/common/directionalmobility.hh +++ b/opm/models/common/directionalmobility.hh @@ -30,9 +30,11 @@ #include #include +#include +#include + #include #include -#include #include namespace Opm { @@ -54,16 +56,15 @@ public: : mobility_{mX, mY, mZ} {} - const array_type& getArray(unsigned index) const + OPM_HOST_DEVICE const array_type& getArray(unsigned index) const { if (index > 2) { - throw std::runtime_error("Unexpected mobility array index " + std::to_string(index)); + OPM_THROW(std::runtime_error, "Unexpected mobility array index"); } - return mobility_[index]; } - array_type& getArray(unsigned index) + OPM_HOST_DEVICE array_type& getArray(unsigned index) { return const_cast(std::as_const(*this).getArray(index)); } diff --git a/opm/models/discretization/common/tpfalinearizer.hh b/opm/models/discretization/common/tpfalinearizer.hh index 94f6267663..333657d7f0 100644 --- a/opm/models/discretization/common/tpfalinearizer.hh +++ b/opm/models/discretization/common/tpfalinearizer.hh @@ -98,7 +98,7 @@ struct FullDomain }; #if HAVE_CUDA && OPM_IS_COMPILING_WITH_GPU_COMPILER - FullDomain> copy_to_gpu(FullDomain<> CPUDomain) + inline FullDomain> copy_to_gpu(FullDomain<> CPUDomain) { if (CPUDomain.cells.size() == 0) { OPM_THROW(std::runtime_error, "Cannot copy empty full domain to GPU."); @@ -108,7 +108,7 @@ struct FullDomain }; }; - FullDomain> make_view(FullDomain>& buffer) + inline FullDomain> make_view(FullDomain>& buffer) { if (buffer.cells.size() == 0) { OPM_THROW(std::runtime_error, "Cannot make view of empty full domain buffer."); diff --git a/opm/simulators/flow/NewTranFluxModule.hpp b/opm/simulators/flow/NewTranFluxModule.hpp index ecfcfa05d7..83811e38cf 100644 --- a/opm/simulators/flow/NewTranFluxModule.hpp +++ b/opm/simulators/flow/NewTranFluxModule.hpp @@ -34,6 +34,7 @@ #include #include +#include #include #include @@ -142,7 +143,7 @@ public: */ OPM_HOST_DEVICE const DimMatrix& intrinsicPermeability() const { - throw std::invalid_argument("The ECL transmissibility module does not provide an explicit intrinsic permeability"); + OPM_THROW(std::invalid_argument, "The ECL transmissibility module does not provide an explicit intrinsic permeability"); } /*! @@ -153,7 +154,7 @@ public: */ OPM_HOST_DEVICE const EvalDimVector& potentialGrad(unsigned) const { - throw std::invalid_argument("The ECL transmissibility module does not provide explicit potential gradients"); + OPM_THROW(std::invalid_argument, "The ECL transmissibility module does not provide explicit potential gradients"); } /*! @@ -173,7 +174,7 @@ public: */ OPM_HOST_DEVICE const EvalDimVector& filterVelocity(unsigned) const { - throw std::invalid_argument("The ECL transmissibility module does not provide explicit filter velocities"); + OPM_THROW(std::invalid_argument, "The ECL transmissibility module does not provide explicit filter velocities"); } /*! diff --git a/opm/simulators/flow/TemperatureModel.hpp b/opm/simulators/flow/TemperatureModel.hpp index f387ecc3d4..704a7cd5c0 100644 --- a/opm/simulators/flow/TemperatureModel.hpp +++ b/opm/simulators/flow/TemperatureModel.hpp @@ -29,6 +29,7 @@ #define OPM_TEMPERATURE_MODEL_HPP #include +#include #include #include @@ -97,15 +98,17 @@ class BlackOilEnergyIntensiveQuantitiesTemp - void updateTemperature_(const Problem& problem, unsigned globalDofIdx, unsigned timeIdx) + OPM_HOST_DEVICE void updateTemperature_(const Problem& problem, + unsigned globalDofIdx, + unsigned timeIdx) { const EvaluationTemp T = EvaluationTemp::createVariable(problem.temperature(globalDofIdx, timeIdx), 0); fluidState_.setTemperature(T); } - void updateEnergyQuantities_(const Problem& problem, - const unsigned globalSpaceIdx, - const unsigned timeIdx) + OPM_HOST_DEVICE void updateEnergyQuantities_(const Problem& problem, + const unsigned globalSpaceIdx, + const unsigned timeIdx) { // compute the specific enthalpy of the fluids, the specific enthalpy of the rock // and the thermal conductivity coefficients @@ -132,20 +135,20 @@ class BlackOilEnergyIntensiveQuantitiesTemp rockFraction_ = problem.rockFraction(globalSpaceIdx, timeIdx); } - const EvaluationTemp& rockInternalEnergy() const + OPM_HOST_DEVICE const EvaluationTemp& rockInternalEnergy() const { return rockInternalEnergy_; } - const EvaluationTemp& totalThermalConductivity() const + OPM_HOST_DEVICE const EvaluationTemp& totalThermalConductivity() const { return totalThermalConductivity_; } - const Scalar& rockFraction() const + OPM_HOST_DEVICE const Scalar& rockFraction() const { return rockFraction_; } - const FluidStateTemp& fluidStateTemp() const + OPM_HOST_DEVICE const FluidStateTemp& fluidStateTemp() const { return fluidState_; } template - void setFluidState(const FluidState& fs) + OPM_HOST_DEVICE void setFluidState(const FluidState& fs) { // copy the needed part of the fluid state fluidState_.setPvtRegionIndex(fs.pvtRegionIndex()); diff --git a/tests/gpuistl/test_gpu_linear_two_phase_material.cu b/tests/gpuistl/test_gpu_linear_two_phase_material.cu index 68512d9ab8..604de549a2 100644 --- a/tests/gpuistl/test_gpu_linear_two_phase_material.cu +++ b/tests/gpuistl/test_gpu_linear_two_phase_material.cu @@ -33,6 +33,7 @@ #include #include #include +#include #include #include @@ -80,8 +81,10 @@ BOOST_AUTO_TEST_CASE(TestSimpleInterpolation) NorneEvaluation* gpuAdInput; NorneEvaluation* gpuAdResOnGPU; - OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdInput, sizeof(NorneEvaluation))); - OPM_GPU_SAFE_CALL(cudaMalloc(&gpuAdResOnGPU, sizeof(NorneEvaluation))); + auto gpuAdInputOwner = Opm::gpuistl::make_gpu_unique_ptr(); + auto gpuAdResOnGPUOwner = Opm::gpuistl::make_gpu_unique_ptr(); + gpuAdInput = gpuAdInputOwner.get(); + gpuAdResOnGPU = gpuAdResOnGPUOwner.get(); for (Scalar x_i : testXs){ auto cpuMadeAd = NorneEvaluation(x_i, 0); @@ -96,7 +99,4 @@ BOOST_AUTO_TEST_CASE(TestSimpleInterpolation) BOOST_CHECK(gpuAdResOnCPU == cpuInterpolatedEval); } - - OPM_GPU_SAFE_CALL(cudaFree(gpuAdInput)); - OPM_GPU_SAFE_CALL(cudaFree(gpuAdResOnGPU)); }