/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  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 "gpuSLTSDdtScheme.H"
#include "gpusurfaceInterpolate.H"
#include "gpufvMatrices.H"

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

namespace Foam
{

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

namespace fv
{

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

template<bool add>
struct SLTSDdtSchemeRelaxedDiag
{
    const scalar* phi;
   
    SLTSDdtSchemeRelaxedDiag
    (
        const scalar* _phi
    ):
        phi(_phi)
    {}

    __host__ __device__
    scalar operator()(const label& cell, const label& face)
    {
        if(phi[face])
        {
            return add?phi[face]:0;
        }
        else
        {
            return (!add)?-phi[face]:0;
        }
    }
};

template<class Type>
void gpuSLTSDdtScheme<Type>::relaxedDiag
(
    scalargpuField& rD,
    const surfaceScalargpuField& phi
) const
{
    scalargpuField diag(rD.size(), 0.0);

    matrixOperation
    (
        thrust::make_constant_iterator(0),
        diag,
        mesh().lduAddr(),
        SLTSDdtSchemeRelaxedDiag<true>(phi.data()),
        SLTSDdtSchemeRelaxedDiag<false>(phi.data())
    );

    matrixOperation
    (
        thrust::make_constant_iterator(scalar(0)),
        rD,
        mesh().lduAddr(),
        SLTSDdtSchemeRelaxedDiag<false>(phi.data()),
        SLTSDdtSchemeRelaxedDiag<true>(phi.data())
    );

    forAll(phi.boundaryField(), patchi)
    {
        const fvsPatchScalargpuField& pphi = phi.boundaryField()[patchi];

        matrixPatchOperation
        (
            patchi, 
            diag,
            mesh().lduAddr(),
            SLTSDdtSchemeRelaxedDiag<true>(pphi.data())
        );

        matrixPatchOperation
        (
            patchi, 
            rD,
            mesh().lduAddr(),
            SLTSDdtSchemeRelaxedDiag<false>(pphi.data())
        );
    }

    rD += (1.0/alpha_ - 2.0)*diag;
}


template<class Type>
tmp<volScalargpuField> gpuSLTSDdtScheme<Type>::SLrDeltaT() const
{
    const surfaceScalargpuField& phi =
        mesh().hostmesh().objectRegistry::template
            lookupObject<surfaceScalargpuField>(phiName_);

    const dimensionedScalar& deltaT = mesh().time().deltaT();

    tmp<volScalargpuField> trDeltaT
    (
        new volScalargpuField
        (
            IOobject
            (
                "rDeltaT",
                phi.instance(),
                mesh().hostmesh()
            ),
            mesh(),
            dimensionedScalar(dimless/dimTime, Zero),
            extrapolatedCalculatedFvPatchScalargpuField::typeName
        )
    );

    volScalargpuField& rDeltaT = trDeltaT.ref();

    relaxedDiag(rDeltaT, phi);

    if (phi.dimensions() == dimensionSet(0, 3, -1, 0, 0))
    {
        rDeltaT.primitiveFieldRef() = max
        (
            rDeltaT.primitiveField()/mesh().V(),
            scalar(1)/deltaT.value()
        );
    }
    else if (phi.dimensions() == dimensionSet(1, 0, -1, 0, 0))
    {
        const volScalargpuField& rho =
            mesh().hostmesh().objectRegistry::template lookupObject<volScalargpuField>
            (
                rhoName_
            ).oldTime();

        rDeltaT.primitiveFieldRef() = max
        (
            rDeltaT.primitiveField()/(rho.primitiveField()*mesh().V()),
            scalar(1)/deltaT.value()
        );
    }
    else
    {
        FatalErrorInFunction
            << "Incorrect dimensions of phi: " << phi.dimensions()
            << abort(FatalError);
    }

    rDeltaT.correctBoundaryConditions();

    return trDeltaT;
}


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

    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().V0()/mesh().V());

        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>>
gpuSLTSDdtScheme<Type>::fvcDdt
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(SLrDeltaT());

    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().V0()/mesh().V()
                ),
                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>>
gpuSLTSDdtScheme<Type>::fvcDdt
(
    const dimensionedScalar& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(SLrDeltaT());

    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().V0()/mesh().V()
                ),
                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>>
gpuSLTSDdtScheme<Type>::fvcDdt
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(SLrDeltaT());

    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().V0()/mesh().V()
                ),
                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>>
gpuSLTSDdtScheme<Type>::fvcDdt
(
    const volScalargpuField& alpha,
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& vf
)
{
    const volScalargpuField rDeltaT(SLrDeltaT());

    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>>
gpuSLTSDdtScheme<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(SLrDeltaT()().primitiveField());

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

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

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuSLTSDdtScheme<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(SLrDeltaT()().primitiveField());

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

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

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuSLTSDdtScheme<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(SLrDeltaT()().primitiveField());

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

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

    return tfvm;
}


template<class Type>
tmp<gpufvMatrix<Type>>
gpuSLTSDdtScheme<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(SLrDeltaT()().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 gpuSLTSDdtScheme<Type>::fluxFieldType>
gpuSLTSDdtScheme<Type>::fvcDdtUfCorr
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& Uf
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(SLrDeltaT()));

    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 gpuSLTSDdtScheme<Type>::fluxFieldType>
gpuSLTSDdtScheme<Type>::fvcDdtPhiCorr
(
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const fluxFieldType& phi
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(SLrDeltaT()));

    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 gpuSLTSDdtScheme<Type>::fluxFieldType>
gpuSLTSDdtScheme<Type>::fvcDdtUfCorr
(
    const volScalargpuField& rho,
    const GeometricgpuField<Type, fvPatchgpuField, gpuvolMesh>& U,
    const GeometricgpuField<Type, fvsPatchgpuField, gpusurfaceMesh>& Uf
)
{
    const surfaceScalargpuField rDeltaT(fvc::interpolate(SLrDeltaT()));

    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 gpuSLTSDdtScheme<Type>::fluxFieldType>
gpuSLTSDdtScheme<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> gpuSLTSDdtScheme<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

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