/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  F ield         | OpenFOAM: The Open Source CFD Toolbox
   \\    /   O peration     |
    \\  /    A nd           | www.openfoam.com
     \\/     M anipulation  |
-------------------------------------------------------------------------------
    Copyright (C) 2011-2018 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 "gpuCoEulerDdtScheme.H"
#include "gpusurfaceInterpolate.H"
//#include "fvcDiv.H"
#include "gpufvMatrices.H"

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

namespace Foam
{

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

namespace fv
{

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

struct CoEulerDdtSchemeCorDeltaTFunctor
{
    const scalar * cofrDeltaT;

    CoEulerDdtSchemeCorDeltaTFunctor
    (
        const scalar * _cofrDeltaT
    ):
        cofrDeltaT(_cofrDeltaT)
    {}
    __host__ __device__
    scalar operator()(const label& cell, const label& face)
    {
        return cofrDeltaT[face];
    }
};
template<class Type>
tmp<volScalargpuField> gpuCoEulerDdtScheme<Type>::CorDeltaT() const
{
    const surfaceScalargpuField cofrDeltaT(CofrDeltaT());

    tmp<volScalargpuField> tcorDeltaT
    (
        new volScalargpuField
        (
            IOobject
            (
                "CorDeltaT",
                cofrDeltaT.instance(),
                mesh().hostmesh()
            ),
            mesh(),
            dimensionedScalar(cofrDeltaT.dimensions(), Zero),
            extrapolatedCalculatedFvPatchScalargpuField::typeName
        )
    );

    volScalargpuField& corDeltaT = tcorDeltaT.ref();
   
    CoEulerDdtSchemeCorDeltaTFunctor fun(cofrDeltaT.data());

    matrixOperation
    (
        thrust::make_constant_iterator(scalar(0)),
        corDeltaT.primitiveFieldRef(),
        mesh().lduAddr(),
        fun,
        fun,
        maxOp<scalar>(),
        maxOp<scalar>()
    );

    const surfaceScalargpuField::Boundary& cofrDeltaTbf =
        cofrDeltaT.boundaryField();

    forAll(cofrDeltaTbf, patchi)
    {
        const fvsPatchScalargpuField& pcofrDeltaT = cofrDeltaTbf[patchi];

        CoEulerDdtSchemeCorDeltaTFunctor pfun(pcofrDeltaT.data());
        matrixPatchOperation
        (
            patchi, 
            corDeltaT.primitiveFieldRef(),
            mesh().lduAddr(),
            pfun,
            maxOp<scalar>()
        );
    }

    corDeltaT.correctBoundaryConditions();

    return tcorDeltaT;
}


template<class Type>
tmp<surfaceScalargpuField> gpuCoEulerDdtScheme<Type>::CofrDeltaT() const
{
    const dimensionedScalar& deltaT = mesh().time().deltaT();

    const surfaceScalargpuField& phi =
        static_cast<const objectRegistry&>(mesh().hostmesh())
        .lookupObject<surfaceScalargpuField>(phiName_);

    if (phi.dimensions() == dimensionSet(0, 3, -1, 0, 0))
    {
        surfaceScalargpuField Co
        (
            mesh().gpusurfaceInterpolation::deltaCoeffs()
           *(mag(phi)/mesh().magSf())
           *deltaT
        );

        return max(Co/maxCo_, scalar(1))/deltaT;
    }
    else if (phi.dimensions() == dimensionSet(1, 0, -1, 0, 0))
    {
        const volScalargpuField& rho =
            static_cast<const objectRegistry&>(mesh().hostmesh())
           .lookupObject<volScalargpuField>(rhoName_).oldTime();

        surfaceScalargpuField Co
        (
            mesh().gpusurfaceInterpolation::deltaCoeffs()
           *(mag(phi)/(fvc::interpolate(rho)*mesh().magSf()))
           *deltaT
        );

        return max(Co/maxCo_, scalar(1))/deltaT;
    }

    FatalErrorInFunction
        << "Incorrect dimensions of phi: " << phi.dimensions()
        << abort(FatalError);

    return nullptr;
}


template<class Type>
tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
gpuCoEulerDdtScheme<Type>::fvcDdt
(
    const dimensioned<Type>& dt
)
{
    const volScalargpuField rDeltaT(CorDeltaT());

    IOobject ddtIOobject
    (
        "ddt("+dt.name()+')',
        mesh().time().timeName(),
        mesh().hostmesh()
    );

    if (mesh().hostmesh().moving())
    {
        tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>> tdtdt
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                dimensioned<Type>(dt.dimensions()/dimTime, Zero)
            )
        );

        tdtdt.ref().primitiveFieldRef() =
            rDeltaT.primitiveField()*dt.value()
           *(1.0 - mesh().Vsc0()/mesh().Vsc());

        return tdtdt;
    }
    else
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                dimensioned<Type>(dt.dimensions()/dimTime, Zero),
                calculatedFvPatchgpuField<Type>::typeName
            )
        );
    }
}


template<class Type>
tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
gpuCoEulerDdtScheme<Type>::fvcDdt
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(CorDeltaT());

    IOobject ddtIOobject
    (
        "ddt("+vf.name()+')',
        mesh().time().timeName(),
        mesh().hostmesh()
    );

    if (mesh().hostmesh().moving())
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                rDeltaT.dimensions()*vf.dimensions(),
                rDeltaT.primitiveField()*
                (
                    vf.primitiveField()
                  - vf.oldTime().primitiveField()*mesh().Vsc0()/mesh().Vsc()
                ),
                rDeltaT.boundaryField()*
                (
                    vf.boundaryField() - vf.oldTime().boundaryField()
                )
            )
        );
    }
    else
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                rDeltaT*(vf - vf.oldTime())
            )
        );
    }
}


template<class Type>
tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
gpuCoEulerDdtScheme<Type>::fvcDdt
(
    const dimensionedScalar& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(CorDeltaT());

    IOobject ddtIOobject
    (
        "ddt("+rho.name()+','+vf.name()+')',
        mesh().time().timeName(),
        mesh().hostmesh()
    );

    if (mesh().hostmesh().moving())
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                rDeltaT.dimensions()*rho.dimensions()*vf.dimensions(),
                rDeltaT.primitiveField()*rho.value()*
                (
                    vf.primitiveField()
                  - vf.oldTime().primitiveField()*mesh().Vsc0()/mesh().Vsc()
                ),
                rDeltaT.boundaryField()*rho.value()*
                (
                    vf.boundaryField() - vf.oldTime().boundaryField()
                )
            )
        );
    }
    else
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                rDeltaT*rho*(vf - vf.oldTime())
            )
        );
    }
}


template<class Type>
tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
gpuCoEulerDdtScheme<Type>::fvcDdt
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(CorDeltaT());

    IOobject ddtIOobject
    (
        "ddt("+rho.name()+','+vf.name()+')',
        mesh().time().timeName(),
        mesh().hostmesh()
    );

    if (mesh().hostmesh().moving())
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                rDeltaT.dimensions()*rho.dimensions()*vf.dimensions(),
                rDeltaT.primitiveField()*
                (
                    rho.primitiveField()*vf.primitiveField()
                  - rho.oldTime().primitiveField()
                   *vf.oldTime().primitiveField()*mesh().Vsc0()/mesh().Vsc()
                ),
                rDeltaT.boundaryField()*
                (
                    rho.boundaryField()*vf.boundaryField()
                  - rho.oldTime().boundaryField()
                   *vf.oldTime().boundaryField()
                )
            )
        );
    }
    else
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                rDeltaT*(rho*vf - rho.oldTime()*vf.oldTime())
            )
        );
    }
}


template<class Type>
tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
gpuCoEulerDdtScheme<Type>::fvcDdt
(
    const volScalargpuField& alpha,
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(CorDeltaT());

    IOobject ddtIOobject
    (
        "ddt("+alpha.name()+','+rho.name()+','+vf.name()+')',
        mesh().time().timeName(),
        mesh().hostmesh()
    );

    if (mesh().hostmesh().moving())
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                mesh(),
                rDeltaT.dimensions()
               *alpha.dimensions()*rho.dimensions()*vf.dimensions(),
                rDeltaT.primitiveField()*
                (
                    alpha.primitiveField()
                   *rho.primitiveField()
                   *vf.primitiveField()

                  - alpha.oldTime().primitiveField()
                   *rho.oldTime().primitiveField()
                   *vf.oldTime().primitiveField()*mesh().Vsc0()/mesh().Vsc()
                ),
                rDeltaT.boundaryField()*
                (
                    alpha.boundaryField()
                   *rho.boundaryField()
                   *vf.boundaryField()

                  - alpha.oldTime().boundaryField()
                   *rho.oldTime().boundaryField()
                   *vf.oldTime().boundaryField()
                )
            )
        );
    }
    else
    {
        return tmp<GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>>
        (
            new GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>
            (
                ddtIOobject,
                rDeltaT
               *(
                   alpha*rho*vf
                 - alpha.oldTime()*rho.oldTime()*vf.oldTime()
                )
            )
        );
    }
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuCoEulerDdtScheme<Type>::fvmDdt
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    tmp<gpufvMatrix<Type>> tfvm
    (
        new gpufvMatrix<Type>
        (
            vf,
            vf.dimensions()*dimVol/dimTime
        )
    );

    gpufvMatrix<Type>& fvm = tfvm.ref();

    scalargpuField rDeltaT(CorDeltaT()().primitiveField());

    fvm.gpuDiag() = rDeltaT*mesh().Vsc();

    if (mesh().hostmesh().moving())
    {
        fvm.source() = rDeltaT*vf.oldTime().primitiveField()*mesh().Vsc0();
    }
    else
    {
        fvm.source() = rDeltaT*vf.oldTime().primitiveField()*mesh().Vsc();
    }

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuCoEulerDdtScheme<Type>::fvmDdt
(
    const dimensionedScalar& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    tmp<gpufvMatrix<Type>> tfvm
    (
        new gpufvMatrix<Type>
        (
            vf,
            rho.dimensions()*vf.dimensions()*dimVol/dimTime
        )
    );
    gpufvMatrix<Type>& fvm = tfvm.ref();

    scalargpuField rDeltaT(CorDeltaT()().primitiveField());

    fvm.gpuDiag() = rDeltaT*rho.value()*mesh().Vsc();

    if (mesh().hostmesh().moving())
    {
        fvm.source() = rDeltaT
            *rho.value()*vf.oldTime().primitiveField()*mesh().Vsc0();
    }
    else
    {
        fvm.source() = rDeltaT
            *rho.value()*vf.oldTime().primitiveField()*mesh().Vsc();
    }

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuCoEulerDdtScheme<Type>::fvmDdt
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    tmp<gpufvMatrix<Type>> tfvm
    (
        new gpufvMatrix<Type>
        (
            vf,
            rho.dimensions()*vf.dimensions()*dimVol/dimTime
        )
    );
    gpufvMatrix<Type>& fvm = tfvm.ref();

    scalargpuField rDeltaT(CorDeltaT()().primitiveField());

    fvm.gpuDiag() = rDeltaT*rho.primitiveField()*mesh().Vsc();

    if (mesh().hostmesh().moving())
    {
        fvm.source() = rDeltaT
            *rho.oldTime().primitiveField()
            *vf.oldTime().primitiveField()*mesh().Vsc0();
    }
    else
    {
        fvm.source() = rDeltaT
            *rho.oldTime().primitiveField()
            *vf.oldTime().primitiveField()*mesh().Vsc();
    }

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuCoEulerDdtScheme<Type>::fvmDdt
(
    const volScalargpuField& alpha,
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    tmp<gpufvMatrix<Type>> tfvm
    (
        new gpufvMatrix<Type>
        (
            vf,
            alpha.dimensions()*rho.dimensions()*vf.dimensions()*dimVol/dimTime
        )
    );
    gpufvMatrix<Type>& fvm = tfvm.ref();

    scalargpuField rDeltaT(CorDeltaT()().primitiveField());

    fvm.gpuDiag() =
        rDeltaT*alpha.primitiveField()*rho.primitiveField()*mesh().Vsc();

    if (mesh().hostmesh().moving())
    {
        fvm.source() = rDeltaT
            *alpha.oldTime().primitiveField()
            *rho.oldTime().primitiveField()
            *vf.oldTime().primitiveField()*mesh().Vsc0();
    }
    else
    {
        fvm.source() = rDeltaT
            *alpha.oldTime().primitiveField()
            *rho.oldTime().primitiveField()
            *vf.oldTime().primitiveField()*mesh().Vsc();
    }

    return tfvm;
}


template<class Type>
tmp<typename gpuCoEulerDdtScheme<Type>::fluxFieldType>
gpuCoEulerDdtScheme<Type>::fvcDdtUfCorr
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& Uf
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(CorDeltaT()));

    fluxFieldType phiUf0(mesh().Sf() & Uf.oldTime());
    fluxFieldType phiCorr
    (
        phiUf0 - fvc::dotInterpolate(mesh().Sf(), U.oldTime())
    );

    return tmp<fluxFieldType>
    (
        new fluxFieldType
        (
            IOobject
            (
                "ddtCorr(" + U.name() + ',' + Uf.name() + ')',
                mesh().time().timeName(),
                mesh().hostmesh()
            ),
            this->fvcDdtPhiCoeff(U.oldTime(), phiUf0, phiCorr)
           *rDeltaT*phiCorr
        )
    );
}


template<class Type>
tmp<typename gpuCoEulerDdtScheme<Type>::fluxFieldType>
gpuCoEulerDdtScheme<Type>::fvcDdtPhiCorr
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const fluxFieldType& phi
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(CorDeltaT()));

    fluxFieldType phiCorr
    (
        phi.oldTime() - fvc::dotInterpolate(mesh().Sf(), U.oldTime())
    );

    return tmp<fluxFieldType>
    (
        new fluxFieldType
        (
            IOobject
            (
                "ddtCorr(" + U.name() + ',' + phi.name() + ')',
                mesh().time().timeName(),
                mesh().hostmesh()
            ),
            this->fvcDdtPhiCoeff(U.oldTime(), phi.oldTime(), phiCorr)
           *rDeltaT*phiCorr
        )
    );
}


template<class Type>
tmp<typename gpuCoEulerDdtScheme<Type>::fluxFieldType>
gpuCoEulerDdtScheme<Type>::fvcDdtUfCorr
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& Uf
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(CorDeltaT()));

    if
    (
        U.dimensions() == dimVelocity
     && Uf.dimensions() == dimDensity*dimVelocity
    )
    {
        GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh> rhoU0
        (
            rho.oldTime()*U.oldTime()
        );

        fluxFieldType phiUf0(mesh().Sf() & Uf.oldTime());
        fluxFieldType phiCorr(phiUf0 - fvc::dotInterpolate(mesh().Sf(), rhoU0));

        return tmp<fluxFieldType>
        (
            new fluxFieldType
            (
                IOobject
                (
                    "ddtCorr("
                  + rho.name() + ',' + U.name() + ',' + Uf.name() + ')',
                    mesh().time().timeName(),
                    mesh().hostmesh()
                ),
                this->fvcDdtPhiCoeff(rhoU0, phiUf0, phiCorr, rho.oldTime())
               *rDeltaT*phiCorr
            )
        );
    }
    else if
    (
        U.dimensions() == dimDensity*dimVelocity
     && Uf.dimensions() == dimDensity*dimVelocity
    )
    {
        fluxFieldType phiUf0(mesh().Sf() & Uf.oldTime());
        fluxFieldType phiCorr
        (
            phiUf0 - fvc::dotInterpolate(mesh().Sf(), U.oldTime())
        );

        return tmp<fluxFieldType>
        (
            new fluxFieldType
            (
                IOobject
                (
                    "ddtCorr("
                  + rho.name() + ',' + U.name() + ',' + Uf.name() + ')',
                    mesh().time().timeName(),
                    mesh().hostmesh()
                ),
                this->fvcDdtPhiCoeff
                (
                    U.oldTime(),
                    phiUf0,
                    phiCorr,
                    rho.oldTime()
                )*rDeltaT*phiCorr
            )
        );
    }
    else
    {
        FatalErrorInFunction
            << "dimensions of Uf are not correct"
            << abort(FatalError);

        return fluxFieldType::null();
    }
}


template<class Type>
tmp<typename gpuCoEulerDdtScheme<Type>::fluxFieldType>
gpuCoEulerDdtScheme<Type>::fvcDdtPhiCorr
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const fluxFieldType& phi
)
{
    dimensionedScalar rDeltaT = 1.0/mesh().time().deltaT();

    if
    (
        U.dimensions() == dimVelocity
     && phi.dimensions() == rho.dimensions()*dimVelocity*dimArea
    )
    {
        GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh> rhoU0
        (
            rho.oldTime()*U.oldTime()
        );

        fluxFieldType phiCorr
        (
            phi.oldTime() - fvc::dotInterpolate(mesh().Sf(), rhoU0)
        );

        return tmp<fluxFieldType>
        (
            new fluxFieldType
            (
                IOobject
                (
                    "ddtCorr("
                  + rho.name() + ',' + U.name() + ',' + phi.name() + ')',
                    mesh().time().timeName(),
                    mesh().hostmesh()
                ),
                this->fvcDdtPhiCoeff
                (
                    rhoU0,
                    phi.oldTime(),
                    phiCorr,
                    rho.oldTime()
                )*rDeltaT*phiCorr
            )
        );
    }
    else if
    (
        U.dimensions() == rho.dimensions()*dimVelocity
     && phi.dimensions() == rho.dimensions()*dimVelocity*dimArea
    )
    {
        fluxFieldType phiCorr
        (
            phi.oldTime() - fvc::dotInterpolate(mesh().Sf(), U.oldTime())
        );

        return tmp<fluxFieldType>
        (
            new fluxFieldType
            (
                IOobject
                (
                    "ddtCorr("
                  + rho.name() + ',' + U.name() + ',' + phi.name() + ')',
                    mesh().time().timeName(),
                    mesh().hostmesh()
                ),
                this->fvcDdtPhiCoeff
                (
                    U.oldTime(),
                    phi.oldTime(),
                    phiCorr,
                    rho.oldTime()
                )*rDeltaT*phiCorr
            )
        );
    }
    else
    {
        FatalErrorInFunction
            << "dimensions of phi are not correct"
            << abort(FatalError);

        return fluxFieldType::null();
    }
}


template<class Type>
tmp<surfaceScalargpuField> gpuCoEulerDdtScheme<Type>::meshPhi
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>&
)
{
    tmp<surfaceScalargpuField> tmeshPhi
    (
        new surfaceScalargpuField
        (
            IOobject
            (
                "meshPhi",
                mesh().time().timeName(),
                mesh().hostmesh(),
                IOobject::NO_READ,
                IOobject::NO_WRITE,
                false
            ),
            mesh(),
            dimensionedScalar(dimVolume/dimTime, Zero)
        )
    );

    tmeshPhi.ref().setOriented();

    return tmeshPhi;
}


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

} // End namespace fv

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

} // End namespace Foam

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