Merge pull request #6676 from multitalentloes/make_newtranfluxmodule_gpu_friendly

make more files gpu-compatible

CMakeLists_files.cmake

@@ -747,6 +747,7 @@ list (APPEND PUBLIC_HEADER_FILES
   opm/models/blackoil/blackoillocalresidualtpfa.hh
   opm/models/blackoil/blackoilmeanings.hh
   opm/models/blackoil/blackoilmodel.hh
+  opm/models/blackoil/blackoilmoduleparams.hh
   opm/models/blackoil/blackoilnewtonmethod.hpp
   opm/models/blackoil/blackoilnewtonmethodparams.hpp
   opm/models/blackoil/blackoilonephaseindices.hh

flowexperimental/BlackOilIntensiveQuantitiesGlobalIndex.hpp

@@ -522,6 +522,14 @@ public:
         throw std::logic_error("permFactor() is not yet implemented for compositional modeling");
     }
 
+    /*!
+     * \brief Returns the fluid system used by this intensive quantities.
+     */
+    OPM_HOST_DEVICE const auto& getFluidSystem() const
+    {
+        return fluidState_.fluidSystem();
+    }
+
 private:
     friend BlackOilSolventIntensiveQuantities<TypeTag, enableSolvent>;
     friend BlackOilExtboIntensiveQuantities<TypeTag, enableExtbo>;

opm/models/blackoil/blackoilconvectivemixingmodule.hh

@@ -40,6 +40,10 @@
 #include <opm/models/blackoil/blackoilenergymodules.hh>
 #include <opm/models/discretization/common/fvbaseproperties.hh>
 
+#include <opm/common/utility/VectorWithDefaultAllocator.hpp>
+
+#include <opm/common/utility/gpuistl_if_available.hpp>
+
 #if HAVE_ECL_INPUT
 #include <cstddef>
 #endif
@@ -48,6 +52,40 @@
 
 namespace Opm {
 
+    template<class Scalar, template<class> class Storage = Opm::VectorWithDefaultAllocator>
+    struct ConvectiveMixingModuleParam
+    {
+        Storage<bool> active_;
+        Storage<Scalar> Xhi_;
+        Storage<Scalar> Psi_;
+    };
+
+#ifdef HAVE_CUDA
+namespace gpuistl
+{
+    template <class Scalar>
+    ConvectiveMixingModuleParam<Scalar, GpuView> make_view(ConvectiveMixingModuleParam<Scalar, GpuBuffer>& params)
+    {
+        ConvectiveMixingModuleParam<Scalar, GpuView> view;
+        view.active_ = gpuistl::make_view(params.active_);
+        view.Xhi_ = gpuistl::make_view(params.Xhi_);
+        view.Psi_ = gpuistl::make_view(params.Psi_);
+        return view;
+    }
+
+    template <class Scalar>
+    ConvectiveMixingModuleParam<Scalar, GpuBuffer> copy_to_gpu(const ConvectiveMixingModuleParam<Scalar, Opm::VectorWithDefaultAllocator>& params)
+    {
+        return ConvectiveMixingModuleParam<Scalar, GpuBuffer>{
+            gpuistl::GpuBuffer(params.active_),
+            gpuistl::GpuBuffer(params.Xhi_),
+            gpuistl::GpuBuffer(params.Psi_)
+        };
+    }
+}
+
+#endif
+
 /*!
  * \copydoc Opm::BlackOilConvectiveMixingModule
  * \brief Provides the convective term in the transport flux for the brine
@@ -77,14 +115,12 @@ class BlackOilConvectiveMixingModule<TypeTag, /*enableConvectiveMixing=*/false>
     using IntensiveQuantities = GetPropType<TypeTag, Properties::IntensiveQuantities>;
 
 public:
-    struct ConvectiveMixingModuleParam
-    {};
 
     #if HAVE_ECL_INPUT
     static void beginEpisode(const EclipseState&,
                              const Schedule&,
                              const int,
-                             ConvectiveMixingModuleParam&)
+                             ConvectiveMixingModuleParam<Scalar>&)
     {}
     #endif
 
@@ -96,7 +132,7 @@ public:
                                  const IntensiveQuantities&,
                                  const IntensiveQuantities&,
                                  const unsigned int,
-                                 const ConvectiveMixingModuleParam&)
+                                 const ConvectiveMixingModuleParam<Scalar>&)
     {}
 
     template <class Context>
@@ -118,7 +154,7 @@ public:
                                         const Scalar,
                                         const Scalar,
                                         const Scalar,
-                                        const ConvectiveMixingModuleParam&)
+                                        const ConvectiveMixingModuleParam<Scalar>&)
     {}
 };
 
@@ -133,6 +169,7 @@ class BlackOilConvectiveMixingModule<TypeTag, /*enableConvectiveMixing=*/true>
     using IntensiveQuantities = GetPropType<TypeTag, Properties::IntensiveQuantities>;
     using GridView = GetPropType<TypeTag, Properties::GridView>;
     using Toolbox = MathToolbox<Evaluation>;
+    using ConvectiveMixingModuleParamT = ConvectiveMixingModuleParam<Scalar>;
 
     enum { conti0EqIdx = Indices::conti0EqIdx };
     enum { dimWorld = GridView::dimensionworld };
@@ -142,18 +179,11 @@ class BlackOilConvectiveMixingModule<TypeTag, /*enableConvectiveMixing=*/true>
     static constexpr unsigned contiEnergyEqIdx = Indices::contiEnergyEqIdx;
 
 public:
-    struct ConvectiveMixingModuleParam
-    {
-        std::vector<bool> active_;
-        std::vector<Scalar> Xhi_;
-        std::vector<Scalar> Psi_;
-    };
-
     #if HAVE_ECL_INPUT
     static void beginEpisode(const EclipseState& eclState,
                              const Schedule& schedule,
                              const int episodeIdx,
-                             ConvectiveMixingModuleParam& info)
+                             ConvectiveMixingModuleParamT& info)
     {
         // check that Xhi and Psi didn't change
         std::size_t numRegions = eclState.runspec().tabdims().getNumPVTTables();
@@ -173,11 +203,12 @@ public:
     }
     #endif
 
-    static void modifyAvgDensity(Evaluation& rhoAvg,
-                                 const IntensiveQuantities& intQuantsIn,
-                                 const IntensiveQuantities& intQuantsEx,
-                                 const unsigned phaseIdx,
-                                 const ConvectiveMixingModuleParam& info) {
+    template <class CMMParam>
+    OPM_HOST_DEVICE static void modifyAvgDensity(Evaluation& rhoAvg,
+                                                 const IntensiveQuantities& intQuantsIn,
+                                                 const IntensiveQuantities& intQuantsEx,
+                                                 const unsigned phaseIdx,
+                                                 const CMMParam& info) {
 
         if (info.active_.empty()) {
             return;
@@ -186,14 +217,15 @@ public:
             return;
         }
 
-        if (phaseIdx == FluidSystem::gasPhaseIdx) {
+        FluidSystem fsys = intQuantsIn.getFluidSystem();
+        if (phaseIdx == fsys.gasPhaseIdx) {
             return;
         }
 
         const auto& liquidPhaseIdx =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-               ? FluidSystem::waterPhaseIdx
-               : FluidSystem::oilPhaseIdx;
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+               ? fsys.waterPhaseIdx
+               : fsys.oilPhaseIdx;
 
         // Compute avg density based on pure water
         const auto& t_in = intQuantsIn.fluidState().temperature(liquidPhaseIdx);
@@ -201,16 +233,16 @@ public:
         const auto& salt_in = intQuantsIn.fluidState().saltConcentration();
 
         const auto& bLiquidIn =
-            FluidSystem::phaseIsActive(waterPhaseIdx)
-                ? FluidSystem::waterPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
+            fsys.phaseIsActive(waterPhaseIdx)
+                ? fsys.waterPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
                                                                          t_in, p_in, Evaluation(0.0), salt_in)
-                : FluidSystem::oilPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
+                : fsys.oilPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
                                                                  t_in, p_in, Evaluation(0.0));
 
         const auto& refDensityLiquidIn =
-            FluidSystem::phaseIsActive(waterPhaseIdx)
-                ? FluidSystem::waterPvt().waterReferenceDensity(intQuantsIn.pvtRegionIndex())
-                : FluidSystem::oilPvt().oilReferenceDensity(intQuantsIn.pvtRegionIndex());
+            fsys.phaseIsActive(waterPhaseIdx)
+                ? fsys.waterPvt().waterReferenceDensity(intQuantsIn.pvtRegionIndex())
+                : fsys.oilPvt().oilReferenceDensity(intQuantsIn.pvtRegionIndex());
 
         const auto& rho_in =  bLiquidIn * refDensityLiquidIn;
 
@@ -219,16 +251,16 @@ public:
         const auto salt_ex = Toolbox::value(intQuantsEx.fluidState().saltConcentration());
 
         const auto bLiquidEx =
-            FluidSystem::phaseIsActive(waterPhaseIdx)
-                ? FluidSystem::waterPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
+            fsys.phaseIsActive(waterPhaseIdx)
+                ? fsys.waterPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
                                                                        t_ex, p_ex, Scalar{0.0}, salt_ex)
-                : FluidSystem::oilPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
+                : fsys.oilPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
                                                                      t_ex, p_ex, Scalar{0.0});
 
         const auto& refDensityLiquidEx =
-            FluidSystem::phaseIsActive(waterPhaseIdx)
-                ? FluidSystem::waterPvt().waterReferenceDensity(intQuantsEx.pvtRegionIndex())
-                : FluidSystem::oilPvt().oilReferenceDensity(intQuantsEx.pvtRegionIndex());
+            fsys.phaseIsActive(waterPhaseIdx)
+                ? fsys.waterPvt().waterReferenceDensity(intQuantsEx.pvtRegionIndex())
+                : fsys.oilPvt().oilReferenceDensity(intQuantsEx.pvtRegionIndex());
 
         const auto rho_ex =  bLiquidEx * refDensityLiquidEx;
 
@@ -236,10 +268,10 @@ public:
     }
 
     template <class Context>
-    static void addConvectiveMixingFlux(RateVector& flux,
-                                        const Context& elemCtx,
-                                        unsigned scvfIdx,
-                                        unsigned timeIdx)
+    OPM_HOST_DEVICE static void addConvectiveMixingFlux(RateVector& flux,
+                                                        const Context& elemCtx,
+                                                        unsigned scvfIdx,
+                                                        unsigned timeIdx)
     {
         // need for darcy flux calculation
         const auto& problem = elemCtx.problem();
@@ -275,16 +307,19 @@ public:
      * \brief Adds the convective mixing mass flux flux to the flux vector over a flux
      *        integration point.
       */
-    static void addConvectiveMixingFlux(RateVector& flux,
-                                        const IntensiveQuantities& intQuantsIn,
-                                        const IntensiveQuantities& intQuantsEx,
-                                        const unsigned globalIndexIn,
-                                        const unsigned globalIndexEx,
-                                        const Scalar distZg,
-                                        const Scalar trans,
-                                        const Scalar faceArea,
-                                        const ConvectiveMixingModuleParam& info)
+    template <class RateVectorT = RateVector, class CMMParam = ConvectiveMixingModuleParamT>
+    OPM_HOST_DEVICE static void addConvectiveMixingFlux(RateVectorT& flux,
+                                                        const IntensiveQuantities& intQuantsIn,
+                                                        const IntensiveQuantities& intQuantsEx,
+                                                        const unsigned globalIndexIn,
+                                                        const unsigned globalIndexEx,
+                                                        const Scalar distZg,
+                                                        const Scalar trans,
+                                                        const Scalar faceArea,
+                                                        const CMMParam& info)
     {
+        const FluidSystem& fsys = intQuantsIn.getFluidSystem();
+
         if (info.active_.empty()) {
             return;
         }
@@ -294,9 +329,9 @@ public:
         }
 
         const auto& liquidPhaseIdx =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-                ? FluidSystem::waterPhaseIdx
-                : FluidSystem::oilPhaseIdx;
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+                ? fsys.waterPhaseIdx
+                : fsys.oilPhaseIdx;
 
         // interiour
         const auto& t_in = intQuantsIn.fluidState().temperature(liquidPhaseIdx);
@@ -305,21 +340,21 @@ public:
         const auto& salt_in = intQuantsIn.fluidState().saltSaturation();
 
         const auto bLiquidSatIn =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-                ? FluidSystem::waterPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+                ? fsys.waterPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
                                                                        t_in, p_in, rssat_in, salt_in)
-                : FluidSystem::oilPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
+                : fsys.oilPvt().inverseFormationVolumeFactor(intQuantsIn.pvtRegionIndex(),
                                                                      t_in, p_in, rssat_in);
 
         const auto& densityLiquidIn =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-                ? FluidSystem::waterPvt().waterReferenceDensity(intQuantsIn.pvtRegionIndex())
-                : FluidSystem::oilPvt().oilReferenceDensity(intQuantsIn.pvtRegionIndex());
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+                ? fsys.waterPvt().waterReferenceDensity(intQuantsIn.pvtRegionIndex())
+                : fsys.oilPvt().oilReferenceDensity(intQuantsIn.pvtRegionIndex());
 
         const auto rho_in = Opm::getValue(intQuantsIn.fluidState().invB(liquidPhaseIdx)) * densityLiquidIn;
         const auto rho_sat_in = bLiquidSatIn *
                                 (densityLiquidIn +
-                                 rssat_in * FluidSystem::referenceDensity(FluidSystem::gasPhaseIdx,
+                                 rssat_in * fsys.referenceDensity(fsys.gasPhaseIdx,
                                                                           intQuantsIn.pvtRegionIndex()));
 
         // exteriour
@@ -328,21 +363,21 @@ public:
         const auto rssat_ex = Opm::getValue(intQuantsEx.saturatedDissolutionFactor());
         const auto salt_ex = Opm::getValue(intQuantsEx.fluidState().saltSaturation());
         const auto bLiquidSatEx =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-                ? FluidSystem::waterPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+                ? fsys.waterPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
                                                                        t_ex, p_ex, rssat_ex, salt_ex)
-                : FluidSystem::oilPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
+                : fsys.oilPvt().inverseFormationVolumeFactor(intQuantsEx.pvtRegionIndex(),
                                                                      t_ex, p_ex, rssat_ex);
 
         const auto& densityLiquidEx =
-            FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
-                ? FluidSystem::waterPvt().waterReferenceDensity(intQuantsEx.pvtRegionIndex())
-                : FluidSystem::oilPvt().oilReferenceDensity(intQuantsEx.pvtRegionIndex());
+            fsys.phaseIsActive(fsys.waterPhaseIdx)
+                ? fsys.waterPvt().waterReferenceDensity(intQuantsEx.pvtRegionIndex())
+                : fsys.oilPvt().oilReferenceDensity(intQuantsEx.pvtRegionIndex());
 
         const auto rho_ex = Opm::getValue(intQuantsEx.fluidState().invB(liquidPhaseIdx)) * densityLiquidEx;
         const auto rho_sat_ex = bLiquidSatEx *
                                 (densityLiquidEx +
-                                 rssat_ex * FluidSystem::referenceDensity(FluidSystem::gasPhaseIdx,
+                                 rssat_ex * fsys.referenceDensity(fsys.gasPhaseIdx,
                                                                           intQuantsEx.pvtRegionIndex()));
         // rho difference approximation
         const auto delta_rho = (rho_sat_ex + rho_sat_in - rho_in - rho_ex) / 2;
@@ -363,7 +398,7 @@ public:
             const auto& rssat_up = (upIdx == interiorDofIdx) ? rssat_in : rssat_ex;
             unsigned globalUpIndex = (upIdx == interiorDofIdx) ? globalIndexIn : globalIndexEx;
             const auto& Rsup =
-                FluidSystem::phaseIsActive(FluidSystem::waterPhaseIdx)
+                fsys.phaseIsActive(fsys.waterPhaseIdx)
                     ? up.fluidState().Rsw()
                     : up.fluidState().Rs();
 
@@ -376,7 +411,7 @@ public:
 
             const auto convectiveFlux = -trans * transMult * info.Xhi_[up.pvtRegionIndex()] * invB *
                                         pressure_difference_convective_mixing * RsupRestricted / (visc * faceArea);
-            unsigned activeGasCompIdx = FluidSystem::canonicalToActiveCompIdx(FluidSystem::gasCompIdx);
+            unsigned activeGasCompIdx = fsys.canonicalToActiveCompIdx(fsys.gasCompIdx);
             if (globalUpIndex == globalIndexIn) {
                 flux[conti0EqIdx + activeGasCompIdx] += convectiveFlux;
             }
@@ -386,7 +421,7 @@ public:
 
             if constexpr (enableFullyImplicitThermal) {
                 const auto& h = up.fluidState().enthalpy(liquidPhaseIdx) *
-                                FluidSystem::referenceDensity(FluidSystem::gasPhaseIdx, up.pvtRegionIndex());
+                                fsys.referenceDensity(fsys.gasPhaseIdx, up.pvtRegionIndex());
                 if (globalUpIndex == globalIndexIn) {
                     flux[contiEnergyEqIdx] += convectiveFlux * h;
                 }
@@ -434,7 +469,7 @@ public:
                                                                               SoMax);
     }
 
-    const Evaluation& saturatedDissolutionFactor() const
+    OPM_HOST_DEVICE const Evaluation& saturatedDissolutionFactor() const
     { return saturatedDissolutionFactor_; }
 
 protected:

opm/models/blackoil/blackoilintensivequantities.hh

@@ -814,6 +814,14 @@ public:
         }
     }
 
+    /*!
+     * \brief Returns the fluid system used by this intensive quantities.
+     */
+    OPM_HOST_DEVICE const auto& getFluidSystem() const
+    {
+        return fluidState_.fluidSystem();
+    }
+
 private:
     friend BlackOilSolventIntensiveQuantities<TypeTag, enableSolvent>;
     friend BlackOilExtboIntensiveQuantities<TypeTag, enableExtbo>;

opm/models/blackoil/blackoillocalresidualtpfa.hh

@@ -43,6 +43,7 @@
 #include <opm/models/blackoil/blackoilenergymodules.hh>
 #include <opm/models/blackoil/blackoilextbomodules.hh>
 #include <opm/models/blackoil/blackoilfoammodules.hh>
+#include <opm/models/blackoil/blackoilmoduleparams.hh>
 #include <opm/models/blackoil/blackoilpolymermodules.hh>
 #include <opm/models/blackoil/blackoilproperties.hh>
 #include <opm/models/blackoil/blackoilsolventmodules.hh>
@@ -117,7 +118,7 @@ class BlackOilLocalResidualTPFA : public GetPropType<TypeTag, Properties::DiscLo
     using BrineModule = BlackOilBrineModule<TypeTag>;
     using DiffusionModule = BlackOilDiffusionModule<TypeTag, enableDiffusion>;
     using ConvectiveMixingModule = BlackOilConvectiveMixingModule<TypeTag, enableConvectiveMixing>;
-    using ConvectiveMixingModuleParam = typename ConvectiveMixingModule::ConvectiveMixingModuleParam;
+    using ModuleParams = BlackoilModuleParams<ConvectiveMixingModuleParam<Scalar>>;
 
     using DispersionModule = BlackOilDispersionModule<TypeTag, enableDispersion>;
     using BioeffectsModule = BlackOilBioeffectsModule<TypeTag>;
@@ -140,11 +141,6 @@ public:
         ConditionalStorage<enableDispersion, double> dispersivity;
     };
 
-    struct ModuleParams
-    {
-        ConvectiveMixingModuleParam convectiveMixingModuleParam;
-    };
-
     /*!
      * \copydoc FvBaseLocalResidual::computeStorage
      */

opm/models/blackoil/blackoilmoduleparams.hh

@@ -0,0 +1,27 @@
+/*
+  Copyright 2025 Equinor ASA
+  This file is part of the Open Porous Media project (OPM).
+  OPM is free software: you can redistribute it and/or modify
+  it under the terms of the GNU General Public License as published by
+  the Free Software Foundation, either version 3 of the License, or
+  (at your option) any later version.
+  OPM is distributed in the hope that it will be useful,
+  but WITHOUT ANY WARRANTY; without even the implied warranty of
+  MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the
+  GNU General Public License for more details.
+  You should have received a copy of the GNU General Public License
+  along with OPM.  If not, see <http://www.gnu.org/licenses/>.
+*/
+
+#ifndef OPM_MODELS_BLACKOIL_MODULEPARAM_HH
+#define OPM_MODELS_BLACKOIL_MODULEPARAM_HH
+
+namespace Opm {
+    template<class ConvectiveMixingModuleParamT>
+    struct BlackoilModuleParams
+    {
+        ConvectiveMixingModuleParamT convectiveMixingModuleParam;
+    };
+} // namespace Opm
+
+#endif // OPM_MODELS_BLACKOIL_MODULEPARAM_HH

opm/simulators/flow/FlowProblemBlackoil.hpp

@@ -39,7 +39,8 @@
 #include <opm/material/fluidsystems/blackoilpvt/ConstantCompressibilityOilPvt.hpp>
 #include <opm/material/fluidsystems/blackoilpvt/ConstantCompressibilityWaterPvt.hpp>
 
-#include <opm/models/blackoil/blackoillocalresidualtpfa.hh>
+#include <opm/models/blackoil/blackoilconvectivemixingmodule.hh>
+#include <opm/models/blackoil/blackoilmoduleparams.hh>
 
 #include <opm/output/eclipse/EclipseIO.hpp>
 
@@ -148,7 +149,7 @@ private:
     using DispersionModule = BlackOilDispersionModule<TypeTag, enableDispersion>;
     using DiffusionModule = BlackOilDiffusionModule<TypeTag, enableDiffusion>;
     using ConvectiveMixingModule = BlackOilConvectiveMixingModule<TypeTag, enableConvectiveMixing>;
-    using ModuleParams = typename BlackOilLocalResidualTPFA<TypeTag>::ModuleParams;
+    using ModuleParams = BlackoilModuleParams<ConvectiveMixingModuleParam<Scalar>>;
     using HybridNewton = BlackOilHybridNewton<TypeTag>;
 
     using InitialFluidState = typename EquilInitializer<TypeTag>::ScalarFluidState;

opm/simulators/flow/NewTranFluxModule.hpp

@@ -44,7 +44,10 @@
 #include <opm/models/discretization/common/fvbaseproperties.hh>
 #include <opm/models/blackoil/blackoilproperties.hh>
 #include <opm/models/utils/signum.hh>
-#include <opm/models/blackoil/blackoillocalresidualtpfa.hh>
+#include <opm/models/blackoil/blackoilmoduleparams.hh>
+#include <opm/models/blackoil/blackoilconvectivemixingmodule.hh>
+
+#include <opm/common/utility/gpuDecorators.hpp>
 
 #include <array>
 
@@ -132,12 +135,12 @@ class NewTranExtensiveQuantities
     using DimMatrix = Dune::FieldMatrix<Scalar, dimWorld, dimWorld>;
 
     using ConvectiveMixingModule = BlackOilConvectiveMixingModule<TypeTag, enableConvectiveMixing>;
-    using ModuleParams = typename BlackOilLocalResidualTPFA<TypeTag>::ModuleParams;
+    using ModuleParams = BlackoilModuleParams<ConvectiveMixingModuleParam<Scalar>>;
 public:
     /*!
      * \brief Return the intrinsic permeability tensor at a face [m^2]
      */
-    const DimMatrix& intrinsicPermeability() const
+    OPM_HOST_DEVICE const DimMatrix& intrinsicPermeability() const
     {
         throw std::invalid_argument("The ECL transmissibility module does not provide an explicit intrinsic permeability");
     }
@@ -148,7 +151,7 @@ public:
      *
      * \param phaseIdx The index of the fluid phase
      */
-    const EvalDimVector& potentialGrad(unsigned) const
+    OPM_HOST_DEVICE const EvalDimVector& potentialGrad(unsigned) const
     {
         throw std::invalid_argument("The ECL transmissibility module does not provide explicit potential gradients");
     }
@@ -159,7 +162,7 @@ public:
      *
      * \param phaseIdx The index of the fluid phase
      */
-    const Evaluation& pressureDifference(unsigned phaseIdx) const
+    OPM_HOST_DEVICE const Evaluation& pressureDifference(unsigned phaseIdx) const
     { return pressureDifference_[phaseIdx]; }
 
     /*!
@@ -168,7 +171,7 @@ public:
      *
      * \param phaseIdx The index of the fluid phase
      */
-    const EvalDimVector& filterVelocity(unsigned) const
+    OPM_HOST_DEVICE const EvalDimVector& filterVelocity(unsigned) const
     {
         throw std::invalid_argument("The ECL transmissibility module does not provide explicit filter velocities");
     }
@@ -182,7 +185,7 @@ public:
      *
      * \param phaseIdx The index of the fluid phase
      */
-    const Evaluation& volumeFlux(unsigned phaseIdx) const
+    OPM_HOST_DEVICE const Evaluation& volumeFlux(unsigned phaseIdx) const
     { return volumeFlux_[phaseIdx]; }
 
 protected:
@@ -193,7 +196,7 @@ protected:
      * i.e., the DOF which exhibits a higher effective pressure for
      * the given phase.
      */
-    unsigned upstreamIndex_(unsigned phaseIdx) const
+    OPM_HOST_DEVICE unsigned upstreamIndex_(unsigned phaseIdx) const
     {
         assert(phaseIdx < numPhases);
 
@@ -207,22 +210,22 @@ protected:
      * i.e., the DOF which exhibits a lower effective pressure for the
      * given phase.
      */
-    unsigned downstreamIndex_(unsigned phaseIdx) const
+    OPM_HOST_DEVICE unsigned downstreamIndex_(unsigned phaseIdx) const
     {
         assert(phaseIdx < numPhases);
 
         return dnIdx_[phaseIdx];
     }
 
-    void updateSolvent(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
+    OPM_HOST_DEVICE void updateSolvent(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
     { asImp_().updateVolumeFluxTrans(elemCtx, scvfIdx, timeIdx); }
 
-    void updatePolymer(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
+    OPM_HOST_DEVICE void updatePolymer(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
     { asImp_().updateShearMultipliers(elemCtx, scvfIdx, timeIdx); }
 
 public:
 
-    static void volumeAndPhasePressureDifferences(std::array<short, numPhases>& upIdx,
+    OPM_HOST_DEVICE static void volumeAndPhasePressureDifferences(std::array<short, numPhases>& upIdx,
                                                   std::array<short, numPhases>& dnIdx,
                                                   Evaluation (&volumeFlux)[numPhases],
                                                   Evaluation (&pressureDifferences)[numPhases],
@@ -307,7 +310,7 @@ public:
     }
 
     template<class EvalType>
-    static void calculatePhasePressureDiff_(short& upIdx,
+    OPM_HOST_DEVICE static void calculatePhasePressureDiff_(short& upIdx,
                                             short& dnIdx,
                                             EvalType& pressureDifference,
                                             const IntensiveQuantities& intQuantsIn,
@@ -414,7 +417,7 @@ protected:
     /*!
      * \brief Update the required gradients for interior faces
      */
-    void calculateGradients_(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
+    OPM_HOST_DEVICE void calculateGradients_(const ElementContext& elemCtx, unsigned scvfIdx, unsigned timeIdx)
     {
         Valgrind::SetUndefined(*this);
 
@@ -425,7 +428,7 @@ protected:
      * \brief Update the required gradients for boundary faces
      */
     template <class FluidState>
-    void calculateBoundaryGradients_(const ElementContext& elemCtx,
+    OPM_HOST_DEVICE void calculateBoundaryGradients_(const ElementContext& elemCtx,
                                      unsigned scvfIdx,
                                      unsigned timeIdx,
                                      const FluidState& exFluidState)
@@ -469,7 +472,7 @@ public:
      * \brief Update the required gradients for boundary faces
      */
     template <class Problem, class FluidState, class EvaluationContainer>
-    static void calculateBoundaryGradients_(const Problem& problem,
+    OPM_HOST_DEVICE static void calculateBoundaryGradients_(const Problem& problem,
                                             const unsigned globalSpaceIdx,
                                             const IntensiveQuantities& intQuantsIn,
                                             const unsigned bfIdx,
@@ -568,10 +571,10 @@ protected:
     /*!
      * \brief Update the volumetric fluxes for all fluid phases on the interior faces of the context
      */
-    void calculateFluxes_(const ElementContext&, unsigned, unsigned)
+    OPM_HOST_DEVICE void calculateFluxes_(const ElementContext&, unsigned, unsigned)
     { }
 
-    void calculateBoundaryFluxes_(const ElementContext&, unsigned, unsigned)
+    OPM_HOST_DEVICE void calculateBoundaryFluxes_(const ElementContext&, unsigned, unsigned)
     {}
 
 private:

opm/simulators/linalg/gpuistl/GpuView.hpp

@@ -121,6 +121,13 @@ public:
         return m_dataPtr;
     }
 
+    /**
+     * @return true if the view has no elements, false otherwise
+     */
+    __host__ __device__ bool empty() const {
+        return m_numberOfElements == 0;
+    }
+
     /**
      * @return fetch the first element in a GpuView
      */