/*---------------------------------------------------------------------------*\
  =========                 |
  \\      /  F ield         | OpenFOAM: The Open Source CFD Toolbox
   \\    /   O peration     |
    \\  /    A nd           | www.openfoam.com
     \\/     M anipulation  |
-------------------------------------------------------------------------------
    Copyright (C) 2011-2014 OpenFOAM Foundation
    Copyright (C) 2019 OpenCFD Ltd.
-------------------------------------------------------------------------------
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 "processorgpuFvPatch.H"
#include "addToRunTimeSelectionTable.H"
#include "transformField.H"

// * * * * * * * * * * * * * * Static Data Members * * * * * * * * * * * * * //

namespace Foam
{
    defineTypeNameAndDebug(processorgpuFvPatch, 0);
    addToRunTimeSelectionTable(gpufvPatch, processorgpuFvPatch, polyPatch);
}


// * * * * * * * * * * * * * * Member Functions  * * * * * * * * * * * * * * //

namespace Foam
{
struct processorFvPatchMakeWeights
{
    template<class Tuple>
    __host__ __device__
    scalar operator()(const Tuple& t)
    {
        const vector& nfa = thrust::get<0>(t);
        const vector& nfc = thrust::get<1>(t);
        const vector& nfcc = thrust::get<2>(t);
        const vector& nf = thrust::get<3>(t);
        const vector& delta = thrust::get<4>(t);

        // The face normals point in the opposite direction on the other side
        vector nfaHat = nfa/(mag(nfa)+VSMALL);
        scalar nfccn = (nfaHat  & (nfc - nfcc));

        return nfccn/((nf&delta)+nfccn);
    }
};
}

void Foam::processorgpuFvPatch::makeWeights(scalargpuField& w) const
{
    if (Pstream::parRun())
    {
        tmp<vectorgpuField> tdelt = coupledgpuFvPatch::delta();
        const vectorgpuField& delt = tdelt();

        tmp<vectorgpuField> tnf = nf();
        const vectorgpuField& nf = tnf();

        auto iter = thrust::make_zip_iterator(thrust::make_tuple(
            procPolyPatch_.gpuNeighbFaceAreas().begin(),
            procPolyPatch_.gpuNeighbFaceCentres().begin(),
            procPolyPatch_.gpuNeighbFaceCellCentres().begin(),
            nf.begin(),
            delt.begin()
        ));
        thrust::transform(iter, iter+size(), w.begin(),
             processorFvPatchMakeWeights()
        );
/*
        // The face normals point in the opposite direction on the other side
        scalarField neighbFaceCentresCn
        (
            (
                procPolyPatch_.neighbFaceAreas()
               /(mag(procPolyPatch_.neighbFaceAreas()) + VSMALL)
            )
          & (
              procPolyPatch_.neighbFaceCentres()
            - procPolyPatch_.neighbFaceCellCentres())
        );

        w = neighbFaceCentresCn
           /((nf()&coupledFvPatch::delta()) + neighbFaceCentresCn);
*/
    }
    else
    {
        w = 1.0;
    }
}


Foam::tmp<Foam::vectorgpuField> Foam::processorgpuFvPatch::delta() const
{
    if (Pstream::parRun())
    {
        // To the transformation if necessary
        if (parallel())
        {
            return
                coupledgpuFvPatch::delta()
              - (
                    procPolyPatch_.gpuNeighbFaceCentres()
                  - procPolyPatch_.gpuNeighbFaceCellCentres()
                );
        }
        else
        {
            return
                coupledgpuFvPatch::delta()
              - transform
                (
                    gpuForwardT(),
                    (
                        procPolyPatch_.gpuNeighbFaceCentres()
                      - procPolyPatch_.gpuNeighbFaceCellCentres()
                    )
                );
        }
    }
    else
    {
        return coupledgpuFvPatch::delta();
    }
}


Foam::tmp<Foam::labelgpuField> Foam::processorgpuFvPatch::interfaceInternalField
(
    const labelgpuList& internalData
) const
{
    return patchInternalField(internalData);
}


Foam::tmp<Foam::labelgpuField> Foam::processorgpuFvPatch::interfaceInternalField
(
    const labelgpuList& internalData,
    const labelgpuList& faceCells
) const
{
    return patchInternalField(internalData, faceCells);
}


void Foam::processorgpuFvPatch::initInternalFieldTransfer
(
    const Pstream::commsTypes commsType,
    const labelgpuList& iF
) const
{
    send(commsType, patchInternalField(iF)());
}


void Foam::processorgpuFvPatch::initInternalFieldTransfer
(
    const Pstream::commsTypes commsType,
    const labelgpuList& iF,
    const labelgpuList& faceCells
) const
{
    send(commsType, interfaceInternalField(iF, faceCells)());
}


Foam::tmp<Foam::labelgpuField> Foam::processorgpuFvPatch::internalFieldTransfer
(
    const Pstream::commsTypes commsType,
    const labelgpuList&
) const
{
    return receive<label>(commsType, this->size());
}


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