/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  F ield         | OpenFOAM: The Open Source CFD Toolbox
   \\    /   O peration     |
    \\  /    A nd           | www.openfoam.com
     \\/     M anipulation  |
-------------------------------------------------------------------------------
    Copyright (C) 2011-2016 OpenFOAM Foundation
-------------------------------------------------------------------------------
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 "gpulinearUpwindV.H"
#include "gpufvMesh.H"
#include "volgpuFields.H"
#include "surfacegpuFields.H"

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

namespace Foam
{

template<class Type, class ProdType>
struct linearUpwindVSfCorrFunctor
{
    const Type zero;

    linearUpwindVSfCorrFunctor
    (
    ):
       zero(pTraits<Type>::zero)
    {}

    template<class Tuple>
    __host__ __device__
    Type operator()(const scalar& faceFlux, const Tuple& t)
    {
        const scalar& w = thrust::get<0>(t);
        const Type& vfn = thrust::get<1>(t);
        const Type& vfo = thrust::get<2>(t);
        const vector& Cf = thrust::get<3>(t);
        const vector& Co = thrust::get<4>(t);
        const vector& Cn = thrust::get<5>(t);
        const ProdType& gradVfo = thrust::get<6>(t);
        const ProdType& gradVfn = thrust::get<7>(t);

        Type maxCorr;
        Type sfCorr;

        if (faceFlux > 0.0)
        {
            maxCorr = (1.0 - w)*(vfn - vfo);

            sfCorr = (Cf - Co) & gradVfo;
        }
        else
        {
            maxCorr = w*(vfo - vfn);

            sfCorr = (Cf - Cn) & gradVfn;
        }

        scalar sfCorrs = magSqr(sfCorr);
        scalar maxCorrs = sfCorr & maxCorr;

        if (sfCorrs > 0)
        {
            if (maxCorrs < 0)
            {
                return zero;
            }
            else if (sfCorrs > maxCorrs)
            {
                return sfCorr * maxCorrs/(sfCorrs + VSMALL);
            }
        }
        else if (sfCorrs < 0)
        {
            if (maxCorrs > 0)
            {
                return zero;
            }
            else if (sfCorrs < maxCorrs)
            {
                return sfCorr * maxCorrs/(sfCorrs - VSMALL);
            }
        }
        return zero;
    }
};


template<class Type, class ProdType>
struct linearUpwindVPatchSfCorrFunctor
{
    const Type zero;

    linearUpwindVPatchSfCorrFunctor
    (
    ):
       zero(pTraits<Type>::zero)
    {}

    template<class Tuple>
    __host__ __device__
    Type operator()(const scalar& pFaceFlux, const Tuple& t)
    {
        const scalar& pW = thrust::get<0>(t);
        const Type& vfo = thrust::get<1>(t);
        const Type& pVfNei = thrust::get<2>(t);
        const vector& pCf = thrust::get<3>(t);
        const vector& Co = thrust::get<4>(t);
        const vector& pd = thrust::get<5>(t);
        const ProdType& gradVfo = thrust::get<6>(t);
        const ProdType& pGradVfNei = thrust::get<7>(t);

        vector maxCorr;
        Type pSfCorr;

        if (pFaceFlux > 0)
        {
            pSfCorr = (pCf - Co) & gradVfo;

            maxCorr = (1.0 - pW)*(pVfNei - vfo);
        }
        else
        {
            pSfCorr = (pCf - pd - Co) & pGradVfNei;

            maxCorr = pW*(vfo - pVfNei);
        }

        scalar pSfCorrs = magSqr(pSfCorr);
        scalar maxCorrs = pSfCorr & maxCorr;

        if (pSfCorrs > 0)
        {
            if (maxCorrs < 0)
            {
                return zero;
            }
            else if (pSfCorrs > maxCorrs)
            {
                return pSfCorr * maxCorrs/(pSfCorrs + VSMALL);
            }
        }
        else if (pSfCorrs < 0)
        {
            if (maxCorrs > 0)
            {
                return zero;
            }
            else if (pSfCorrs < maxCorrs)
            {
                return pSfCorr * maxCorrs/(pSfCorrs - VSMALL);
            }
        }
        return zero;
    }
};

}


template<class Type>
Foam::tmp<Foam::GeometricgpuField<Type, Foam::fvsPatchgpuField, Foam::gpusurfaceMesh>>
Foam::gpulinearUpwindV<Type>::correction
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
) const
{
    const gpufvMesh& mesh = this->mesh();

    tmp<GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>> tsfCorr
    (
        new GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>
        (
            IOobject
            (
                "gpulinearUpwindV::correction(" + vf.name() + ')',
                mesh.time().timeName(),
                mesh.hostmesh(),
                IOobject::NO_READ,
                IOobject::NO_WRITE,
                false
            ),
            mesh,
            dimensioned<Type>(vf.dimensions(), Zero)
        )
    );

    GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& sfCorr = tsfCorr.ref();

    const surfaceScalargpuField& faceFlux = this->faceFlux_;
    const surfaceScalargpuField& w = mesh.weights();

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

    const volVectorgpuField& C = mesh.C();
    const surfaceVectorgpuField& Cf = mesh.Cf();

    tmp
    <
        GeometricgpuField
        <
            typename outerProduct<vector, Type>::type,
            fvPatchgpuField,
            gpuvolMesh
        >
    > tgradVf = gradScheme_().grad(vf, gradSchemeName_);

    const GeometricgpuField
    <
        typename outerProduct<vector, Type>::type,
        fvPatchgpuField,
        gpuvolMesh
    >& gradVf = tgradVf();

    /*forAll(faceFlux, facei)
    {
        vector maxCorr;

        if (faceFlux[facei] > 0.0)
        {
            maxCorr =
                (1.0 - w[facei])*(vf[nei[facei]] - vf[own[facei]]);

            sfCorr[facei] =
                (Cf[facei] - C[own[facei]]) & gradVf[own[facei]];
        }
        else
        {
            maxCorr =
                w[facei]*(vf[own[facei]] - vf[nei[facei]]);

            sfCorr[facei] =
               (Cf[facei] - C[nei[facei]]) & gradVf[nei[facei]];
        }

        scalar sfCorrs = magSqr(sfCorr[facei]);
        scalar maxCorrs = sfCorr[facei] & maxCorr;

        if (sfCorrs > 0)
        {
            if (maxCorrs < 0)
            {
                sfCorr[facei] = Zero;
            }
            else if (sfCorrs > maxCorrs)
            {
                sfCorr[facei] *= maxCorrs/(sfCorrs + VSMALL);
            }
        }
    }*/

    thrust::transform
    (
        faceFlux.field().begin(),
        faceFlux.field().end(),
        thrust::make_zip_iterator(thrust::make_tuple
        (
            w.field().begin(), 
            thrust::make_permutation_iterator
            (
                vf.field().begin(),
                nei.begin()
            ),
            thrust::make_permutation_iterator
            (
                vf.field().begin(),
                own.begin()
            ),
            Cf.field().begin(),
            thrust::make_permutation_iterator
            (
                C.field().begin(),
                own.begin()
            ),
            thrust::make_permutation_iterator
            (
                C.field().begin(),
                nei.begin()
            ),
            thrust::make_permutation_iterator
            (
                gradVf.field().begin(),
                own.begin()
            ),
            thrust::make_permutation_iterator
            (
                gradVf.field().begin(),
                nei.begin()
            )
        )),
        sfCorr.field().begin(),
        linearUpwindVSfCorrFunctor
        <
            Type,
            typename outerProduct<vector, Type>::type
        >()
    );


    typename GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>::
        Boundary& bSfCorr = sfCorr.boundaryFieldRef();

    forAll(bSfCorr, patchi)
    {
        fvsPatchgpuField<Type>& pSfCorr = bSfCorr[patchi];

        if (pSfCorr.coupled())
        {
            const labelgpuList& pOwner =
                mesh.boundary()[patchi].gpuFaceCells();

            const vectorgpuField& pCf = Cf.boundaryField()[patchi];
            const scalargpuField& pW = w.boundaryField()[patchi];

            const scalargpuField& pFaceFlux = faceFlux.boundaryField()[patchi];

            const gpuField<typename outerProduct<vector, Type>::type> pGradVfNei
            (
                gradVf.boundaryField()[patchi].patchNeighbourField()
            );

            const gpuField<Type> pVfNei
            (
                vf.boundaryField()[patchi].patchNeighbourField()
            );

            // Build the d-vectors
            vectorgpuField pd(Cf.boundaryField()[patchi].patch().delta());

            /*forAll(pOwner, facei)
            {
                label own = pOwner[facei];

                vector maxCorr;

                if (pFaceFlux[facei] > 0)
                {
                    pSfCorr[facei] = (pCf[facei] - C[own]) & gradVf[own];

                    maxCorr = (1.0 - pW[facei])*(pVfNei[facei] - vf[own]);
                }
                else
                {
                    pSfCorr[facei] =
                        (pCf[facei] - pd[facei] - C[own]) & pGradVfNei[facei];

                    maxCorr = pW[facei]*(vf[own] - pVfNei[facei]);
                }

                scalar pSfCorrs = magSqr(pSfCorr[facei]);
                scalar maxCorrs = pSfCorr[facei] & maxCorr;

                if (pSfCorrs > 0)
                {
                    if (maxCorrs < 0)
                    {
                        pSfCorr[facei] = Zero;
                    }
                    else if (pSfCorrs > maxCorrs)
                    {
                        pSfCorr[facei] *= maxCorrs/(pSfCorrs + VSMALL);
                    }
                }
            }*/
            thrust::transform
            (
                pFaceFlux.begin(),
                pFaceFlux.end(),
                thrust::make_zip_iterator(thrust::make_tuple
                (
                    pW.begin(),
                    thrust::make_permutation_iterator
                    (
                        vf.field().begin(),
                        pOwner.begin()
                    ),
                    pVfNei.begin(),
                    pCf.begin(),
                    thrust::make_permutation_iterator
                    (
                        C.field().begin(),
                        pOwner.begin()
                    ),
                    pd.begin(),
                    thrust::make_permutation_iterator
                    (
                        gradVf.field().begin(),
                        pOwner.begin()
                    ),
                    pGradVfNei.begin()
                )),
                pSfCorr.begin(), 
                linearUpwindVPatchSfCorrFunctor
                <
                    Type,
                    typename outerProduct<vector, Type>::type
                >()
            );
        }
    }

    return tsfCorr;
}


namespace Foam
{
    makegpulimitedSurfaceInterpolationTypeScheme(gpulinearUpwindV, vector)
}


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