SimpleFIBlackOilModel.hpp

Go to the documentation of this file.

// -*- mode: C++; tab-width: 4; indent-tabs-mode: nil; c-basic-offset: 4 -*-
// vi: set et ts=4 sw=4 sts=4:
/*
  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 2 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/>.
 
  Consult the COPYING file in the top-level source directory of this
  module for the precise wording of the license and the list of
  copyright holders.
*/
#ifndef SIMPLE_FI_BLACK_OIL_MODEL_HPP
#define SIMPLE_FI_BLACK_OIL_MODEL_HPP
 
#include <opm/models/blackoil/blackoilconvectivemixingmodule.hh>
#include <opm/models/blackoil/blackoilmodel.hh>
#include <opm/models/blackoil/blackoilmoduleparams.hh>
#include <opm/models/parallel/threadmanager.hpp>
#include <opm/models/utils/propertysystem.hh>
 
#include <opm/common/ErrorMacros.hpp>
 
#include <opm/grid/CpGrid.hpp>
#include <opm/grid/utility/ElementChunks.hpp>
 
#include <opm/simulators/utils/DeferredLoggingErrorHelpers.hpp>
 
#include <opm/material/fluidmatrixinteractions/EclMultiplexerMaterialParams.hpp>
 
#include <cassert>
#include <cstddef>
#include <stdexcept>
#include <type_traits>
#include <vector>
 
#ifdef _OPENMP
#include <omp.h>
#endif
 
/*
    This file implements a simplified version of the FIBlackOilModel
    which can be used in a GPU environment.
    This is developed while parallelizing the tpfa matrix assembly, so
    only functionality needed there is implemented.
    To start with this means just being able to access intensiveQuantities.
*/
 
namespace Opm
{
 
template <typename TypeTag, template <class> class Storage = Opm::VectorWithDefaultAllocator>
class SimpleFIBlackOilModel
{
public:
    using FluidSystem = GetPropType<TypeTag, Properties::FluidSystem>;
    using TypeTagPublic = TypeTag;
    using Scalar = GetPropType<TypeTag, Properties::Scalar>;
 
    SimpleFIBlackOilModel(
        const Storage<BlackOilIntensiveQuantities<TypeTag>>& cachedIntensiveQuantities0,
        const Storage<BlackOilIntensiveQuantities<TypeTag>>& cachedIntensiveQuantities1,
        const Storage<Scalar>& volumes)
        : cachedIntensiveQuantities0_(cachedIntensiveQuantities0)
        , cachedIntensiveQuantities1_(cachedIntensiveQuantities1)
        , volumes_(volumes)
    {
    }
 
    // Special ctor because fvbasediscretization stores BOIQs with a custom allocator.
    // OtherIQContainer can be any container (e.g. std::vector with a non-default allocator).
    // Could this be using refs instead of making the copy? Then the copy-to-gpu would have to be
    // rewritten a bit
    template <typename OtherIQContainer>
    SimpleFIBlackOilModel(const OtherIQContainer& cachedIntensiveQuantities0,
                          const OtherIQContainer& cachedIntensiveQuantities1,
                          const Storage<Scalar>& volumes)
        : cachedIntensiveQuantities0_(cachedIntensiveQuantities0.begin(),
                                      cachedIntensiveQuantities0.end())
        , cachedIntensiveQuantities1_(cachedIntensiveQuantities1.begin(),
                                      cachedIntensiveQuantities1.end())
        , volumes_(volumes)
    {
    }
 
    SimpleFIBlackOilModel() = delete;
 
    OPM_HOST_DEVICE Scalar dofTotalVolume(unsigned int globalIdx) const
    {
        return volumes_[globalIdx];
    }
 
    OPM_HOST_DEVICE const auto& intensiveQuantities(unsigned int globalIdx,
                                                    unsigned int timeIdx) const
    {
        assert(timeIdx == 0
               || timeIdx
                   == 1); // TODO: do not use assert for this because it can be run in GPU kernel...
        if (timeIdx == 0) {
            return cachedIntensiveQuantities0_[globalIdx];
        } else {
            return cachedIntensiveQuantities1_[globalIdx];
        }
    }
 
public:
    // assumes --enable-storage-cache=false so that we need not worry about updating cpu-caches from
    // gpu kernels
    Storage<BlackOilIntensiveQuantities<TypeTag>> cachedIntensiveQuantities0_;
    Storage<BlackOilIntensiveQuantities<TypeTag>> cachedIntensiveQuantities1_;
    Storage<Scalar> volumes_;
};
 
#if HAVE_CUDA && OPM_IS_COMPILING_WITH_GPU_COMPILER
namespace gpuistl
{
    // For now we make a copy of the simplified CPU model because we need to set the fluid system
    // pointer without breaking copy semantics
    template <class TypeTag, typename GpuFluidSystem>
    auto copy_to_gpu(const SimpleFIBlackOilModel<TypeTag>& cpuModel, const GpuFluidSystem& fsys)
    {
        using Scalar = typename SimpleFIBlackOilModel<TypeTag>::Scalar;
        // In this copy_to_gpu we need to take care of two things:
        // 1) Copy the cachedIntensiveQuantities_ to GPU (go from vector to GpuBuffer to allocate
        // things on GPU) 2) Ensure that the FluidSystem pointer inside each IntensiveQuantities
        // points to a GPU FluidSystem
        //    The pointer should be set before we move the entire thing to the GPU.
 
        // set pointers
        using CorrectTypeTagView =
            typename ::Opm::Properties::TTag::to_gpu_type_t<TypeTag, GpuView>;
        using CorrectBOIQ = BlackOilIntensiveQuantities<CorrectTypeTagView>;
        using SimplifiedGpuBufferModel = SimpleFIBlackOilModel<CorrectTypeTagView, GpuBuffer>;
 
        assert(cpuModel.cachedIntensiveQuantities0_.size() == cpuModel.cachedIntensiveQuantities1_.size());
        const std::size_t nCells = cpuModel.cachedIntensiveQuantities0_.size();
        assert(nCells > 0);
        // CorrectBOIQ is not default-constructible (non-static FluidSystem requires a pointer at
        // construction time). To enable parallel fill, we create a prototype from the first element
        // and use it to pre-size the vector with valid objects, then overwrite in parallel.
        // Assumption: all cells share the same fsys pointer, so the prototype object is a valid
        // stand-in until it is overwritten.
        std::vector<CorrectBOIQ> cpuIntQuantsWithGpuPtr0;
        std::vector<CorrectBOIQ> cpuIntQuantsWithGpuPtr1;
        CorrectBOIQ proto0 = cpuModel.cachedIntensiveQuantities0_[0]
                                    .template withOtherFluidSystem<CorrectTypeTagView>(fsys);
        CorrectBOIQ proto1 = cpuModel.cachedIntensiveQuantities1_[0]
                                    .template withOtherFluidSystem<CorrectTypeTagView>(fsys);
        cpuIntQuantsWithGpuPtr0.resize(nCells, proto0);
        cpuIntQuantsWithGpuPtr1.resize(nCells, proto1);
 
#ifdef _OPENMP
#pragma omp parallel for schedule(static)
#endif
        for (std::size_t i = 1; i < nCells; ++i) {
            cpuIntQuantsWithGpuPtr0[i]
                = cpuModel.cachedIntensiveQuantities0_[i]
                        .template withOtherFluidSystem<CorrectTypeTagView>(fsys);
            cpuIntQuantsWithGpuPtr1[i]
                = cpuModel.cachedIntensiveQuantities1_[i]
                        .template withOtherFluidSystem<CorrectTypeTagView>(fsys);
        }
 
        GpuBuffer<CorrectBOIQ> gpuBuf0(cpuIntQuantsWithGpuPtr0);
        GpuBuffer<CorrectBOIQ> gpuBuf1(cpuIntQuantsWithGpuPtr1);
        GpuBuffer<Scalar> gpuVolumes(cpuModel.volumes_);
 
        // create new blackoil model providing a gpubuffer allocated version of the intQuants
        auto result = SimplifiedGpuBufferModel(
            std::move(gpuBuf0), std::move(gpuBuf1), std::move(gpuVolumes));
 
        return result;
    }
 
    template <typename LocalTypeTag>
    auto make_view(SimpleFIBlackOilModel<LocalTypeTag, GpuBuffer>& gpuSimpleFIBlackOilModel)
    {
        return SimpleFIBlackOilModel<LocalTypeTag, gpuistl::GpuView>(
            make_view(gpuSimpleFIBlackOilModel.cachedIntensiveQuantities0_),
            make_view(gpuSimpleFIBlackOilModel.cachedIntensiveQuantities1_),
            make_view(gpuSimpleFIBlackOilModel.volumes_));
    }
} // namespace gpuistl
#endif
 
} // namespace Opm
 
#endif // SIMPLE_FI_BLACK_OIL_MODEL_HPP