Merge pull request #7027 from kjetilly/intensive_quantities_2

GPU portability: OPM_HOST_DEVICE/OPM_THROW additions and fluidSystem() use

flowexperimental/BlackOilEnergyIntensiveQuantitiesGlobalIndex.hpp

@@ -59,20 +59,20 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex<TypeTag, EnergyModules::Fully
     static constexpr unsigned numPhases = FluidSystem::numPhases;
 
 public:
-    void updateTemperature_([[maybe_unused]] const Problem& problem,
-                            const PrimaryVariables& priVars,
-                            [[maybe_unused]] unsigned globalSpaceIndex,
-                            unsigned timeIdx)
+    OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem,
+                                            const PrimaryVariables& priVars,
+                                            [[maybe_unused]] unsigned globalSpaceIndex,
+                                            unsigned timeIdx)
     {
         auto& fs = Parent::asImp_().fluidState_;
         // set temperature
         fs.setTemperature(priVars.makeEvaluation(temperatureIdx, timeIdx));
     }
-    void updateEnergyQuantities_(const Problem& problem,
-                                 [[maybe_unused]] const PrimaryVariables& priVars,
-                                 unsigned globalSpaceIndex,
-                                 unsigned timeIdx,
-                                 const ParamCache& paramCache)
+    OPM_HOST_DEVICE void updateEnergyQuantities_(const Problem& problem,
+                                                 [[maybe_unused]] const PrimaryVariables& priVars,
+                                                 unsigned globalSpaceIndex,
+                                                 unsigned timeIdx,
+                                                 const ParamCache& paramCache)
     {
         auto& fs = Parent::asImp_().fluidState_;
 
@@ -120,25 +120,27 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex<TypeTag, EnergyModules::Seque
     static constexpr unsigned numPhases = FluidSystem::numPhases;
 public:
 
-    void updateTemperature_([[maybe_unused]] const Problem& problem,
-                            [[maybe_unused]] const PrimaryVariables& priVars,
-                            [[maybe_unused]] unsigned globalSpaceIndex,
-                            [[maybe_unused]] unsigned timeIdx)
+    OPM_HOST_DEVICE void updateTemperature_([[maybe_unused]] const Problem& problem,
+                                            [[maybe_unused]] const PrimaryVariables& priVars,
+                                            [[maybe_unused]] unsigned globalSpaceIndex,
+                                            [[maybe_unused]] unsigned timeIdx)
     {
-        throw std::logic_error("updateTemperature not implemented "
-                        "SequentialImplicitThermal can not be used with"
-                        "global intensive quantites yet.");
+        OPM_THROW(std::logic_error,
+                  "updateTemperature not implemented "
+                  "SequentialImplicitThermal can not be used with"
+                  "global intensive quantites yet.");
     }
 
-    void updateEnergyQuantities_([[maybe_unused]] const Problem& problem,
-                                 [[maybe_unused]] const PrimaryVariables& priVars,
-                                 [[maybe_unused]] unsigned globalSpaceIndex,
-                                 [[maybe_unused]] unsigned timeIdx,
-                                 [[maybe_unused]] const ParamCache& paramCache)
+    OPM_HOST_DEVICE void updateEnergyQuantities_([[maybe_unused]] const Problem& problem,
+                                                 [[maybe_unused]] const PrimaryVariables& priVars,
+                                                 [[maybe_unused]] unsigned globalSpaceIndex,
+                                                 [[maybe_unused]] unsigned timeIdx,
+                                                 [[maybe_unused]] const ParamCache& paramCache)
     {
-        throw std::logic_error("updateEnergyQuantities_ not implemented "
-                "SequentialImplicitThermal can not be used with"
-                "global intensive quantites yet.");
+        OPM_THROW(std::logic_error,
+                  "updateEnergyQuantities_ not implemented "
+                  "SequentialImplicitThermal can not be used with"
+                  "global intensive quantites yet.");
     }
 
 };
@@ -155,21 +157,21 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex<TypeTag, EnergyModules::Const
     using Scalar = GetPropType<TypeTag, Properties::Scalar>;
 
 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<Evaluation>&)
+    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<Evaluation>&)
     { }
 };
 
@@ -185,17 +187,17 @@ class BlackOilEnergyIntensiveQuantitiesGlobalIndex<TypeTag, EnergyModules::NoTem
     using Scalar = GetPropType<TypeTag, Properties::Scalar>;
 
 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<Evaluation>&)
+    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<Evaluation>&)
     { }
 };
 

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<Scalar>(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<Implementation*>(this); }
 
     Evaluation saltConcentration_;
@@ -487,27 +487,27 @@ class BlackOilBrineIntensiveQuantities<TypeTag, false>
     using Scalar = GetPropType<TypeTag, Properties::Scalar>;
 
 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

opm/models/blackoil/blackoilenergymodules.hh

@@ -30,6 +30,7 @@
 
 #include <dune/common/fvector.hh>
 
+#include <opm/common/ErrorMacros.hpp>
 #include <opm/common/utility/gpuDecorators.hpp>
 
 #include <opm/material/common/Tabulated1DFunction.hpp>
@@ -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<Scalar>(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 <class UpstreamEval>
-    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 <class FluidState>
-    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<Implementation*>(this); }
 
     Evaluation rockInternalEnergy_;
@@ -484,51 +485,54 @@ class BlackOilEnergyIntensiveQuantities<TypeTag, EnergyModules::ConstantTemperat
 
 public:
 
-    void updateTemperature_(const ElementContext& elemCtx,
-                            unsigned dofIdx,
-                            unsigned timeIdx)
+    OPM_HOST_DEVICE void updateTemperature_(const ElementContext& elemCtx,
+                                            unsigned dofIdx,
+                                            unsigned timeIdx)
     {
         updateTemperature_(elemCtx.problem(), elemCtx.globalSpaceIndex(dofIdx, timeIdx), timeIdx);
     }
 
     template<class Problem>
-    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<Evaluation>&)
+    OPM_HOST_DEVICE void updateEnergyQuantities_(const ElementContext&,
+                                                 unsigned,
+                                                 unsigned,
+                                                 const typename FluidSystem::template ParameterCache<Evaluation>&)
     {}
 
-    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<Implementation*>(this); }
 };
 
@@ -548,26 +552,28 @@ class BlackOilEnergyIntensiveQuantities<TypeTag, EnergyModules::SequentialImplic
 
 public:
 
-    void updateTemperature_(const Problem& problem, unsigned globalDofIdx, unsigned timeIdx)
+    OPM_HOST_DEVICE void updateTemperature_(const Problem& problem,
+                                            unsigned globalDofIdx,
+                                            unsigned timeIdx)
     {
         // update the temperature for output (without derivatives)
         auto& fs = asImp_().fluidState_;
         fs.setTemperature(problem.temperature(globalDofIdx, timeIdx));
     }
 
-    void updateTemperature_(const ElementContext& elemCtx,
-                            unsigned dofIdx,
-                            unsigned timeIdx)
+    OPM_HOST_DEVICE void updateTemperature_(const ElementContext& elemCtx,
+                                            unsigned dofIdx,
+                                            unsigned timeIdx)
     {
         updateTemperature_(elemCtx.problem(), elemCtx.globalSpaceIndex(dofIdx, timeIdx), timeIdx);
     }
 
     template<class Problem>
-    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<Implementation*>(this); }
 
 };
@@ -617,41 +625,43 @@ class BlackOilEnergyIntensiveQuantities<TypeTag, EnergyModules::NoTemperature>
     using Problem = GetPropType<TypeTag, Properties::Problem>;
 
 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<class Problem>
-    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<Evaluation>&)
+    OPM_HOST_DEVICE void updateEnergyQuantities_(const ElementContext&,
+                                                 unsigned,
+                                                 unsigned,
+                                                 const typename FluidSystem::template ParameterCache<Evaluation>&)
     {}
 
-    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<Implementation*>(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 <class TypeTag>
@@ -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 <class TypeTag>
@@ -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

opm/models/blackoil/blackoilfoammodules.hh

@@ -30,6 +30,7 @@
 
 #include <dune/common/fvector.hh>
 
+#include <opm/common/ErrorMacros.hpp>
 #include <opm/common/OpmLog/OpmLog.hpp>
 #include <opm/common/utility/gpuDecorators.hpp>
 
@@ -184,7 +185,7 @@ public:
                                      Toolbox::template decay<LhsEval>(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.

opm/models/blackoil/blackoilprimaryvariables.hh

@@ -1024,10 +1024,10 @@ private:
 
     template <class Container>
     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 = SimpleModularFluidState<Scalar,
                                                           numPhases,
@@ -1065,7 +1065,7 @@ private:
         #endif
     }
 
-    void setScaledPressure_(Scalar pressure)
+    OPM_HOST_DEVICE void setScaledPressure_(Scalar pressure)
     { (*this)[Indices::pressureSwitchIdx] = pressure / (this->getPressureScale()); }
 
     // NOTE: When adding new member variables, be sure to update the template copy constructor,

opm/models/common/directionalmobility.hh

@@ -30,9 +30,11 @@
 #include <opm/models/discretization/common/fvbaseproperties.hh>
 #include <opm/models/utils/propertysystem.hh>
 
+#include <opm/common/ErrorMacros.hpp>
+#include <opm/common/utility/gpuDecorators.hpp>
+
 #include <array>
 #include <stdexcept>
-#include <string>
 #include <utility>
 
 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<array_type&>(std::as_const(*this).getArray(index));
     }

opm/models/discretization/common/tpfalinearizer.hh

@@ -98,7 +98,7 @@ struct FullDomain
 };
 
 #if HAVE_CUDA && OPM_IS_COMPILING_WITH_GPU_COMPILER
-    FullDomain<gpuistl::GpuBuffer<int>> copy_to_gpu(FullDomain<> CPUDomain)
+    inline FullDomain<gpuistl::GpuBuffer<int>> 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<gpuistl::GpuView<int>> make_view(FullDomain<gpuistl::GpuBuffer<int>>& buffer)
+    inline FullDomain<gpuistl::GpuView<int>> make_view(FullDomain<gpuistl::GpuBuffer<int>>& buffer)
     {
         if (buffer.cells.size() == 0) {
             OPM_THROW(std::runtime_error, "Cannot make view of empty full domain buffer.");

opm/simulators/flow/NewTranFluxModule.hpp

@@ -34,6 +34,7 @@
 #include <dune/common/fvector.hh>
 #include <dune/common/fmatrix.hh>
 
+#include <opm/common/ErrorMacros.hpp>
 #include <opm/common/OpmLog/OpmLog.hpp>
 
 #include <opm/input/eclipse/EclipseState/Grid/FaceDir.hpp>
@@ -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");
     }
 
     /*!

opm/simulators/flow/TemperatureModel.hpp

@@ -29,6 +29,7 @@
 #define OPM_TEMPERATURE_MODEL_HPP
 
 #include <opm/common/OpmLog/OpmLog.hpp>
+#include <opm/common/utility/gpuDecorators.hpp>
 
 #include <opm/models/parallel/gridcommhandles.hh>
 #include <opm/models/utils/propertysystem.hh>
@@ -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 <class FluidState>
-    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());

tests/gpuistl/test_gpu_linear_two_phase_material.cu

@@ -33,6 +33,7 @@
 #include <opm/simulators/linalg/gpuistl/GpuView.hpp>
 #include <opm/simulators/linalg/gpuistl/GpuBuffer.hpp>
 #include <opm/simulators/linalg/gpuistl/detail/gpu_safe_call.hpp>
+#include <opm/simulators/linalg/gpuistl/gpu_smart_pointer.hpp>
 
 #include <array>
 #include <algorithm>
@@ -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<NorneEvaluation>();
+    auto gpuAdResOnGPUOwner = Opm::gpuistl::make_gpu_unique_ptr<NorneEvaluation>();
+    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));
 }