/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  F ield         | OpenFOAM: The Open Source CFD Toolbox
   \\    /   O peration     |
   \\  /    A nd           | www.openfoam.com
     \\/     M anipulation  |
-------------------------------------------------------------------------------
    Copyright (C) 2011-2017 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/>.

Class
    Foam::outletStabilised

Group
    grpFvSurfaceInterpolationSchemes

Description
    Outlet-stabilised interpolation scheme which applies upwind differencing
    to the faces of the cells adjacent to outlets.

    This is particularly useful to stabilise the velocity at entrainment
    boundaries for LES cases using linear or other centred differencing
    schemes.

SourceFiles
    outletStabilised.C

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

#ifndef gpuoutletStabilised_H
#define gpuoutletStabilised_H

#include "gpusurfaceInterpolationScheme.H"
#include "gpuskewCorrectionVectors.H"
#include "gpulinear.H"
#include "gpugaussGrad.H"
#include "zeroGradientFvPatchgpuField.H"
#include "mixedFvPatchgpuField.H"
#include "directionMixedFvPatchgpuField.H"

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

namespace Foam
{

struct outletStabilisedWeightsFunctor
{
    const label nInternalFaces;
    const label* faces;
    const scalar* faceFlux;
    scalar* w;
   
    outletStabilisedWeightsFunctor
    (
        const label _nif,
        const label* _faces,
        const scalar* _faceFlux,
        scalar* _w
    ):
        nInternalFaces(_nif),
        faces(_faces),
        faceFlux(_faceFlux),
        w(_w)
    {}

    __host__ __device__
    void operator()(const cellData& cd)
    {
        for(label i = 0; i < cd.nFaces(); i++)
        {
            label facei = faces[i+cd.getStart()];
            if(facei<nInternalFaces)
            {
                w[facei] = pos(faceFlux[facei]);
            }
        }
    }
};

template<class Type>
struct outletStabilisedCorrectionFunctor
{
    const Type zero;
    const label nInternalFaces;
    const label* faces;
    Type* corr;
   
    outletStabilisedCorrectionFunctor
    (
        const label _nif,
        const label* _faces,
        Type* _corr
    ):
        zero(pTraits<Type>::zero),
        nInternalFaces(_nif),
        faces(_faces),
        corr(_corr)
    {}

    __host__ __device__
    void operator()(const cellData& cd)
    {
        for(label i = 0; i < cd.nFaces(); i++)
        {
            label facei = faces[i+cd.getStart()];
            if(facei<nInternalFaces)
            {
                corr[facei] = zero;
            }
        }
    }
};

/*---------------------------------------------------------------------------*\
                      Class outletStabilised Declaration
\*---------------------------------------------------------------------------*/

template<class Type>
class gpuoutletStabilised
:
    public gpusurfaceInterpolationScheme<Type>
{
    // Private member data

        const surfaceScalargpuField& faceFlux_;
        tmp<gpusurfaceInterpolationScheme<Type>> tScheme_;


    // Private Member Functions

        //- No copy construct
        gpuoutletStabilised(const gpuoutletStabilised&) = delete;

        //- No copy assignment
        void operator=(const gpuoutletStabilised&) = delete;


public:

    //- Runtime type information
    TypeName("outletStabilised");


    // Constructors

        //- Construct from mesh and Istream
        gpuoutletStabilised
        (
            const gpufvMesh& mesh,
            Istream& is
        )
        :
            gpusurfaceInterpolationScheme<Type>(mesh),
            faceFlux_
            (
                mesh.hostmesh().lookupObject<surfaceScalargpuField>
                (
                    word(is)
                )
            ),
            tScheme_
            (
                gpusurfaceInterpolationScheme<Type>::New(mesh, faceFlux_, is)
            )
        {}


        //- Construct from mesh, faceFlux and Istream
        gpuoutletStabilised
        (
            const gpufvMesh& mesh,
            const surfaceScalargpuField& faceFlux,
            Istream& is
        )
        :
            gpusurfaceInterpolationScheme<Type>(mesh),
            faceFlux_(faceFlux),
            tScheme_
            (
                gpusurfaceInterpolationScheme<Type>::New(mesh, faceFlux, is)
            )
        {}


    // Member Functions

        //- Return the interpolation weighting factors
        tmp<surfaceScalargpuField> weights
        (
            const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
        ) const
        {
            tmp<surfaceScalargpuField> tw = tScheme_().weights(vf);
            surfaceScalargpuField& w = tw.ref();

            const gpufvMesh& mesh_ = this->mesh();
            //const cellList& cells = mesh_.cells();
            const cellDatagpuList& cells = mesh_.hostmesh().gpuCells();
            const labelgpuList& faces = mesh_.hostmesh().gpuCellFaces();
            const label nInternalFaces = mesh_.hostmesh().nInternalFaces();

            forAll(vf.boundaryField(), patchi)
            {
                if
                (
                    isA<zeroGradientFvPatchgpuField<Type>>
                        (vf.boundaryField()[patchi])
                 || isA<mixedFvPatchgpuField<Type>>(vf.boundaryField()[patchi])
                 || isA<directionMixedFvPatchgpuField<Type>>
                    (vf.boundaryField()[patchi])
                )
                {
                    const labelgpuList& pFaceCells =
                        mesh_.boundary()[patchi].gpuFaceCells();

                    thrust::for_each
                    (
                        thrust::make_permutation_iterator
                        (
                            cells.begin(),
                            pFaceCells.begin()
                        ),
                        thrust::make_permutation_iterator
                        (
                            cells.begin(),
                            pFaceCells.end()
                        ),
                        outletStabilisedWeightsFunctor
                        (
                            nInternalFaces,
                            faces.data(),
                            faceFlux_.field().data(),
                            w.field().data()
                        )
                    );

                }
            }

            return tw;
        }

        //- Return true if this scheme uses an explicit correction
        virtual bool corrected() const
        {
            return tScheme_().corrected();
        }

        //- Return the explicit correction to the face-interpolate
        //  set to zero on the near-boundary faces where upwind is applied
        virtual tmp<GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>>
        correction
        (
            const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
        ) const
        {
            if (tScheme_().corrected())
            {
                tmp<GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>> tcorr =
                    tScheme_().correction(vf);

                GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& corr =
                    tcorr.ref();

                const gpufvMesh& mesh_ = this->mesh();
                //const cellList& cells = mesh_.cells();
                const cellDatagpuList& cells = mesh_.hostmesh().gpuCells();
                const labelgpuList& faces = mesh_.hostmesh().gpuCellFaces();
                const label nInternalFaces = mesh_.hostmesh().nInternalFaces();

                forAll(vf.boundaryField(), patchi)
                {
                    if
                    (
                        isA<zeroGradientFvPatchgpuField<Type>>
                            (vf.boundaryField()[patchi])
                     || isA<mixedFvPatchgpuField<Type>>
                            (vf.boundaryField()[patchi])
                    )
                    {
                        const labelgpuList& pFaceCells =
                            mesh_.boundary()[patchi].gpuFaceCells();

                        thrust::for_each
                        (
                            thrust::make_permutation_iterator
                            (
                                cells.begin(),
                                pFaceCells.begin()
                            ),
                            thrust::make_permutation_iterator
                            (
                                cells.begin(),
                                pFaceCells.end()
                            ),
                            outletStabilisedCorrectionFunctor<Type>
                            (
                                nInternalFaces,
                                faces.data(),
                                corr.field().data()
                            )
                        );
                    }
                }

                return tcorr;
            }
            else
            {
                return tmp<GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>>
                (
                    nullptr
                );
            }
        }
};


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

} // End namespace Foam

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

#endif

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