/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  F ield         | OpenFOAM: The Open Source CFD Toolbox
   \\    /   O peration     |
    \\  /    A nd           | www.openfoam.com
     \\/     M anipulation  |
-------------------------------------------------------------------------------
    Copyright (C) 2011-2016 OpenFOAM Foundation
    Copyright (C) 2020 OpenCFD Ltd.
-------------------------------------------------------------------------------
License
    This file is part of OpenFOAM.

    OpenFOAM 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.

    OpenFOAM 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 OpenFOAM.  If not, see <http://www.gnu.org/licenses/>.

\*---------------------------------------------------------------------------*/

#include "volgpuFields.H"
#include "surfacegpuFields.H"
#include "gpufvcGrad.H"
#include "coupledFvPatchgpuFields.H"

// * * * * * * * * * * * * Private Member Functions  * * * * * * * * * * * * //

namespace Foam
{
    template<class Limiter,class phiType, class gradPhiType>
    struct LimitedSchemeCalcLimiterFunctor
    {
        const Limiter limiter;

        LimitedSchemeCalcLimiterFunctor(const Limiter& _limiter):limiter(_limiter){}

        template<class Tuple>
        __host__ __device__
        scalar operator()(const scalar& w, const Tuple& t)
	{
            return limiter.limiter
            (
                w,
                thrust::get<0>(t),
                thrust::get<1>(t),
                thrust::get<2>(t),
                thrust::get<3>(t),
                thrust::get<4>(t),
                thrust::get<5>(t) - thrust::get<6>(t)
            );
        }
    };
}

template<class Type, class Limiter, template<class> class LimitFunc>
void Foam::gpuLimitedScheme<Type, Limiter, LimitFunc>::calcLimiter
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& phi,
    surfaceScalargpuField& limiterField
) const
{
    typedef GeometricgpuField<typename Limiter::phiType, fvPatchgpuField, gpuvolMesh>
        VolFieldType;

    typedef GeometricgpuField<typename Limiter::gradPhiType, fvPatchgpuField, gpuvolMesh>
        GradVolFieldType;

    const gpufvMesh& mesh = this->mesh();

    tmp<VolFieldType> tlPhi = LimitFunc<Type>()(phi);
    const VolFieldType& lPhi = tlPhi();

    tmp<GradVolFieldType> tgradc(fvc::grad(lPhi));
    const GradVolFieldType& gradc = tgradc();

    const surfaceScalargpuField& CDweights = mesh.gpusurfaceInterpolation::weights();

    const labelgpuList& owner = mesh.owner();
    const labelgpuList& neighbour = mesh.neighbour();

    const vectorgpuField& C = mesh.C();

    scalargpuField& pLim = limiterField.primitiveFieldRef();

    thrust::transform
    (
        CDweights.field().begin(),
        CDweights.field().end(),
        thrust::make_zip_iterator(thrust::make_tuple
        (
            this->faceFlux_.field().begin(),
            thrust::make_permutation_iterator
            (
                lPhi.field().begin(),
                owner.begin()
            ),
            thrust::make_permutation_iterator
            (
                lPhi.field().begin(),
                neighbour.begin()
            ),
            thrust::make_permutation_iterator
            (
                gradc.field().begin(),
                owner.begin()
            ),
            thrust::make_permutation_iterator
            (
                gradc.field().begin(),
                neighbour.begin()
            ),
            thrust::make_permutation_iterator
            (
                C.begin(),
                neighbour.begin()
            ),
            thrust::make_permutation_iterator
            (
                C.begin(),
                owner.begin()
            )
        )),
        pLim.begin(),
        LimitedSchemeCalcLimiterFunctor
        <
            Limiter,
            typename Limiter::phiType,
            typename Limiter::gradPhiType
        >
        (
            static_cast<const Limiter&>(*this)
        )
    );

    surfaceScalargpuField::Boundary& bLim = limiterField.boundaryFieldRef();

    forAll(bLim, patchi)
    {
        scalargpuField pLim(bLim[patchi]);

        if (bLim[patchi].coupled())
        {
            const scalargpuField& pCDweights = CDweights.boundaryField()[patchi];
            const scalargpuField& pFaceFlux = this->faceFlux_.boundaryField()[patchi];

            const gpuField<typename Limiter::phiType> plPhiP
            (
                lPhi.boundaryField()[patchi].patchInternalField()
            );
            const gpuField<typename Limiter::phiType> plPhiN
            (
                lPhi.boundaryField()[patchi].patchNeighbourField()
            );
            const gpuField<typename Limiter::gradPhiType> pGradcP
            (
                gradc.boundaryField()[patchi].patchInternalField()
            );
            const gpuField<typename Limiter::gradPhiType> pGradcN
            (
                gradc.boundaryField()[patchi].patchNeighbourField()
            );

            // Build the d-vectors
            vectorgpuField pd(CDweights.boundaryField()[patchi].patch().delta());
               
            thrust::transform
            (
                pCDweights.begin(),
                pCDweights.end(),
                thrust::make_zip_iterator(thrust::make_tuple
                (
                    pFaceFlux.begin(),
                    plPhiP.begin(),
                    plPhiN.begin(),
                    pGradcP.begin(),
                    pGradcN.begin(),
                    pd.begin(),
                    thrust::make_constant_iterator(vector(0,0,0))
                )),
                pLim.begin(),
                LimitedSchemeCalcLimiterFunctor
                <
                    Limiter,
                    typename Limiter::phiType,
                    typename Limiter::gradPhiType
                >
                (
                    static_cast<const Limiter&>(*this)
                )
            );
        }
        else
        {
            pLim = 1.0;
        }
    }

    limiterField.setOriented();
}


// * * * * * * * * * * * * Public Member Functions  * * * * * * * * * * * * //

template<class Type, class Limiter, template<class> class LimitFunc>
Foam::tmp<Foam::surfaceScalargpuField>
Foam::gpuLimitedScheme<Type, Limiter, LimitFunc>::limiter
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& phi
) const
{
    const gpufvMesh& mesh = this->mesh();

    const word limiterFieldName(type() + "Limiter(" + phi.name() + ')');

    if (this->mesh().hostmesh().cache("limiter"))
    {
        if (!mesh.hostmesh().foundObject<surfaceScalargpuField>(limiterFieldName))
        {
            surfaceScalargpuField* limiterField
            (
                new surfaceScalargpuField
                (
                    IOobject
                    (
                        limiterFieldName,
                        mesh.time().timeName(),
                        mesh.hostmesh(),
                        IOobject::NO_READ,
                        IOobject::NO_WRITE
                    ),
                    mesh,
                    dimless
                )
            );

            mesh.hostmesh().objectRegistry::store(limiterField);
        }

        surfaceScalargpuField& limiterField =
            mesh.hostmesh().lookupObjectRef<surfaceScalargpuField>(limiterFieldName);

        calcLimiter(phi, limiterField);

        return tmp<surfaceScalargpuField>::New
        (
            limiterFieldName,
            limiterField
        );
    }
    else
    {
        tmp<surfaceScalargpuField> tlimiterField
        (
            new surfaceScalargpuField
            (
                IOobject
                (
                    limiterFieldName,
                    mesh.time().timeName(),
                    mesh.hostmesh()
                ),
                mesh,
                dimless
            )
        );

        calcLimiter(phi, tlimiterField.ref());

        return tlimiterField;
    }
}


// ************************************************************************* //
