/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  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 "gpulinearUpwind.H"
#include "gpufvMesh.H"

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

namespace Foam
{
template<class Type>
struct linearUpwindCorrectionFunctor: std::binary_function<label,thrust::tuple<scalar,vector>,Type>{
	const label* own;
	const label* nei;
	const vector* C;
	const vector* gradf;
	linearUpwindCorrectionFunctor(const label* _own,const label* _nei, const vector* _C, const vector* _gradf):
	                              own(_own),nei(_nei),C(_C),gradf(_gradf){}
    __host__ __device__
    Type operator()(const label& id,const thrust::tuple<scalar,vector>& t){
        label celli = thrust::get<0>(t) > 0?own[id]:nei[id];
        return (thrust::get<1>(t) - C[celli]) & gradf[celli];
    }
};

template<class Type>
struct linearUpwindPatchCorrectionFunctor{
    __host__ __device__
    Type operator()(const thrust::tuple<scalar,vector,vector,vector,vector,vector>& t){
        if(thrust::get<0>(t) > 0)
        {
             return (thrust::get<1>(t) - thrust::get<2>(t)) & thrust::get<4>(t);
        }
        else
        {
             return (thrust::get<1>(t) - thrust::get<2>(t) - thrust::get<3>(t)) & thrust::get<5>(t);
        }
    }
};
}

template<class Type>
Foam::tmp<Foam::GeometricgpuField<Type, Foam::fvsPatchgpuField, Foam::gpusurfaceMesh>>
Foam::gpulinearUpwind<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
            (
                "gpulinearUpwind::correction(" + vf.name() + ')',
                mesh.time().timeName(),
                mesh.hostmesh(),
                IOobject::NO_READ,
                IOobject::NO_WRITE,
                false
            ),
            mesh,
            dimensioned<Type>(vf.name(), vf.dimensions(), Zero)
        )
    );

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

    const surfaceScalargpuField& faceFlux = this->faceFlux_;

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

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

    tmp<fv::gpugradScheme<scalar>> gradScheme_
    (
        fv::gpugradScheme<scalar>::New
        (
            mesh,
            mesh.hostmesh().gradScheme(gradSchemeName_)
        )
    );

    for (direction cmpt = 0; cmpt < pTraits<Type>::nComponents; cmpt++)
    {
        tmp<volVectorgpuField> tgradVf =
            gradScheme_().grad(vf.component(cmpt), gradSchemeName_);

        const volVectorgpuField& gradVf = tgradVf();

        /*forAll(faceFlux, facei)
        {
            const label celli =
                (faceFlux[facei] > 0) ? owner[facei] : neighbour[facei];

            setComponent(sfCorr[facei], cmpt) =
                (Cf[facei] - C[celli]) & gradVf[celli];
        }*/

        thrust::transform(thrust::make_counting_iterator(0),thrust::make_counting_iterator(0)+faceFlux.size(),
                      thrust::make_zip_iterator(thrust::make_tuple(faceFlux.field().begin(),
                                                                   Cf.field().begin()
                                                                  )
                                               ),
                     sfCorr.field().component(cmpt).ref().begin(),
                     linearUpwindCorrectionFunctor<Type>(
                                                   owner.data(),
                                                   neighbour.data(),
                                                   C.field().data(),
                                                   gradVf.field().data()
                                                   ));

        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& pFaceFlux = faceFlux.boundaryField()[patchi];

                const vectorgpuField pGradVfNei
                (
                    gradVf.boundaryField()[patchi].patchNeighbourField()
                );

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

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

                    if (pFaceFlux[facei] > 0)
                    {
                        setComponent(pSfCorr[facei], cmpt) =
                            (pCf[facei] - C[own])
                          & gradVf[own];
                    }
                    else
                    {
                        setComponent(pSfCorr[facei], cmpt) =
                            (pCf[facei] - pd[facei] - C[own])
                          & pGradVfNei[facei];
                    }
                }*/

                thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(pFaceFlux.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()
                                                                  )
                                               ),
                      thrust::make_zip_iterator(thrust::make_tuple(pFaceFlux.end(),
                                                                   pCf.end(),
                                                                   thrust::make_permutation_iterator(C.field().begin(),pOwner.end()),
                                                                   pd.end(),
                                                                   thrust::make_permutation_iterator(gradVf.field().end(),pOwner.begin()),
                                                                   pGradVfNei.end()
                                                                  )
                                               ),
                     pSfCorr.component(cmpt).ref().begin(),
                     linearUpwindPatchCorrectionFunctor<Type>());
            }
        }
    }

    return tsfCorr;
}


namespace Foam
{
struct linearUpwindCorrectionCpFunctor: std::binary_function<label,thrust::tuple<scalar,vector>,vector>{
	const label* own;
	const label* nei;
	const vector* C;
	const tensor* gradf;
	linearUpwindCorrectionCpFunctor(const label* _own,const label* _nei, const vector* _C, const tensor* _gradf):
	                              own(_own),nei(_nei),C(_C),gradf(_gradf){}
    __host__ __device__
    vector operator()(const label& id,const thrust::tuple<scalar,vector>& t){
        label celli = thrust::get<0>(t) > 0?own[id]:nei[id];
        return (thrust::get<1>(t) - C[celli]) & gradf[celli];
    }
};

struct linearUpwindPatchCorrectionCpFunctor{
    __host__ __device__
    vector operator()(const thrust::tuple<scalar,vector,vector,vector,tensor,tensor>& t){
        if(thrust::get<0>(t) > 0)
        {
             return (thrust::get<1>(t) - thrust::get<2>(t)) & thrust::get<4>(t);
        }
        else
        {
             return (thrust::get<1>(t) - thrust::get<2>(t) - thrust::get<3>(t)) & thrust::get<5>(t);
        }
    }
};
}

template<>
Foam::tmp<Foam::surfaceVectorgpuField>
Foam::gpulinearUpwind<Foam::vector>::correction
(
    const volVectorgpuField& vf
) const
{
    const gpufvMesh& mesh = this->mesh();

    tmp<surfaceVectorgpuField> tsfCorr
    (
        new surfaceVectorgpuField
        (
            IOobject
            (
                "gpulinearUpwind::correction(" + vf.name() + ')',
                mesh.time().timeName(),
                mesh.hostmesh(),
                IOobject::NO_READ,
                IOobject::NO_WRITE,
                false
            ),
            mesh,
            dimensioned<vector>(vf.name(), vf.dimensions(), Zero)
        )
    );

    surfaceVectorgpuField& sfCorr = tsfCorr.ref();

    const surfaceScalargpuField& faceFlux = this->faceFlux_;

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

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

    tmp<fv::gpugradScheme<vector>> gradScheme_
    (
        fv::gpugradScheme<vector>::New
        (
            mesh,
            mesh.hostmesh().gradScheme(gradSchemeName_)
        )
    );

    tmp<volTensorgpuField> tgradVf = gradScheme_().grad(vf, gradSchemeName_);
    const volTensorgpuField& gradVf = tgradVf();

    /*forAll(faceFlux, facei)
    {
        const label celli =
            (faceFlux[facei] > 0) ? owner[facei] : neighbour[facei];
        sfCorr[facei] = (Cf[facei] - C[celli]) & gradVf[celli];
    }*/

    /*thrust::transform(thrust::make_counting_iterator(0),thrust::make_counting_iterator(0)+faceFlux.size(),
                      thrust::make_zip_iterator(thrust::make_tuple(faceFlux.getField().begin(),
                                                                   Cf.getField().begin()
                                                                  )
                                               ),
                     sfCorr.getField().begin(),
                     linearUpwindCorrectionCpFunctor(
                                                   owner.data(),
                                                   neighbour.data(),
                                                   C.getField().data(),
                                                   gradVf.getField().data()
                                                   ));

    */
    typename surfaceVectorgpuField::Boundary& bSfCorr = sfCorr.boundaryFieldRef();

    forAll(bSfCorr, patchi)
    {
        fvsPatchVectorgpuField& pSfCorr = bSfCorr[patchi];

        if (pSfCorr.coupled())
        {
            const labelgpuList& pOwner = mesh.boundary()[patchi].gpuFaceCells();
            const vectorgpuField& pCf = Cf.boundaryField()[patchi];
            const scalargpuField& pFaceFlux = faceFlux.boundaryField()[patchi];

            const tensorgpuField pGradVfNei
            (
                gradVf.boundaryField()[patchi].patchNeighbourField()
            );

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

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

                if (pFaceFlux[facei] > 0)
                {
                    pSfCorr[facei] = (pCf[facei] - C[own]) & gradVf[own];
                }
                else
                {
                    pSfCorr[facei] =
                        (pCf[facei] - pd[facei] - C[own]) & pGradVfNei[facei];
                }
            }*/

            thrust::transform(thrust::make_zip_iterator(thrust::make_tuple(pFaceFlux.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()
                                                                  )
                                               ),
                      thrust::make_zip_iterator(thrust::make_tuple(pFaceFlux.end(),
                                                                   pCf.end(),
                                                                   thrust::make_permutation_iterator(C.field().begin(),pOwner.end()),
                                                                   pd.end(),
                                                                   thrust::make_permutation_iterator(gradVf.field().end(),pOwner.begin()),
                                                                   pGradVfNei.end()
                                                                  )
                                               ),
                     pSfCorr.begin(),
                     linearUpwindPatchCorrectionCpFunctor());

        }
    }

    return tsfCorr;
}


namespace Foam
{
    //makelimitedSurfaceInterpolationScheme(linearUpwind)
    makegpulimitedSurfaceInterpolationTypeScheme(gpulinearUpwind, scalar)
    makegpulimitedSurfaceInterpolationTypeScheme(gpulinearUpwind, vector)
}

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