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

// * * * * * * * * * * * * * * * * Constructors  * * * * * * * * * * * * * * //

template<class Type>
Foam::jumpCyclicFvPatchgpuField<Type>::jumpCyclicFvPatchgpuField
(
    const gpufvPatch& p,
    const DimensionedgpuField<Type, gpuvolMesh>& iF
)
:
    cyclicFvPatchgpuField<Type>(p, iF)
{}


template<class Type>
Foam::jumpCyclicFvPatchgpuField<Type>::jumpCyclicFvPatchgpuField
(
    const jumpCyclicFvPatchgpuField<Type>& ptf,
    const gpufvPatch& p,
    const DimensionedgpuField<Type, gpuvolMesh>& iF,
    const fvPatchgpuFieldMapper& mapper
)
:
    cyclicFvPatchgpuField<Type>(ptf, p, iF, mapper)
{}


template<class Type>
Foam::jumpCyclicFvPatchgpuField<Type>::jumpCyclicFvPatchgpuField
(
    const gpufvPatch& p,
    const DimensionedgpuField<Type, gpuvolMesh>& iF,
    const dictionary& dict,
    const bool valueRequired
)
:
    cyclicFvPatchgpuField<Type>(p, iF, dict, valueRequired)
{
    // Call this evaluation in derived classes
    //this->evaluate(Pstream::commsTypes::blocking);
}


template<class Type>
Foam::jumpCyclicFvPatchgpuField<Type>::jumpCyclicFvPatchgpuField
(
    const jumpCyclicFvPatchgpuField<Type>& ptf
)
:
    cyclicFvPatchgpuField<Type>(ptf)
{}


template<class Type>
Foam::jumpCyclicFvPatchgpuField<Type>::jumpCyclicFvPatchgpuField
(
    const jumpCyclicFvPatchgpuField<Type>& ptf,
    const DimensionedgpuField<Type, gpuvolMesh>& iF
)
:
    cyclicFvPatchgpuField<Type>(ptf, iF)
{}


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

namespace Foam
{
    template<class Type>
    struct jumpCyclicTransformFunctor
    {
        const tensor t;
        jumpCyclicTransformFunctor(tensor _t): t(_t){}
        __host__ __device__
        Type operator()(const Type& iF, const Type& jF)
        {
            return transform(t,iF) - jF;
        }
    };
}

template<class Type>
Foam::tmp<Foam::gpuField<Type>>
Foam::jumpCyclicFvPatchgpuField<Type>::patchNeighbourField() const
{
    const gpuField<Type>& iField = this->primitiveField();
    const labelgpuList& nbrFaceCells =
        this->cyclicPatch().neighbFvPatch().gpuFaceCells();

    auto tpnf = tmp<gpuField<Type>>::New(this->size());
    auto& pnf = tpnf.ref();

    gpuField<Type> jf(this->jump());
    if (!this->cyclicPatch().owner())
    {
        jf *= -1.0;
    }

    auto nbrInternalValuesStart = thrust::make_permutation_iterator(
        iField.begin(),nbrFaceCells.begin());

    if (this->doTransform())
    {
        tensor t = this->gpuForwardT().first();
        
        thrust::transform
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + this->size(),
            jf.begin(),
            pnf.begin(),
            jumpCyclicTransformFunctor<Type>(t)
        );
        /*
        forAll(*this, facei)
        {
            pnf[facei] = transform
            (
                this->forwardT()[0], iField[nbrFaceCells[facei]]
            ) - jf[facei];
        }
        */
    }
    else
    {
        thrust::transform
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + this->size(),
            jf.begin(),
            pnf.begin(),
            subtractOperatorFunctor<Type,Type,Type>()
        );
        /*
        forAll(*this, facei)
        {
            pnf[facei] = iField[nbrFaceCells[facei]] - jf[facei];
        }
        */
    }

    return tpnf;
}


template<class Type>
void Foam::jumpCyclicFvPatchgpuField<Type>::updateInterfaceMatrix
(
    scalargpuField& result,
    const bool add,
    const gpulduAddressing& lduAddr,
    const label patchId,
    const scalargpuField& psiInternal,
    const scalargpuField& coeffs,
    const direction cmpt,
    const Pstream::commsTypes
) const
{
    NotImplemented;
}

template<class Type>
void Foam::jumpCyclicFvPatchgpuField<Type>::updateInterfaceMatrix
(
    gpuField<Type>& result,
    const bool add,
    const gpulduAddressing& lduAddr,
    const label patchId,
    const gpuField<Type>& psiInternal,
    const scalargpuField& coeffs,
    const Pstream::commsTypes
) const
{
    gpuField<Type> pnf(this->size());

    const labelgpuList& nbrFaceCells =
        lduAddr.patchAddr
        (
            this->cyclicPatch().neighbPatchID()
        );
    auto nbrInternalValuesStart = thrust::make_permutation_iterator(
        psiInternal.begin(),nbrFaceCells.begin());

    // only apply jump to original field
    if (&psiInternal == &this->primitiveField())
    {
        gpuField<Type> jf(this->jump().ref());

        if (!this->cyclicPatch().owner())
        {
            jf *= -1.0;
        }
		
        thrust::transform
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + this->size(),
            jf.begin(),
            pnf.begin(),
            subtractOperatorFunctor<Type,Type,Type>()
        );
    }
    else
    {
        thrust::copy
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + this->size(),
            pnf.begin()
         );
    }

    // Transform according to the transformation tensors
    this->transformCoupleField(pnf);

    // Multiply the field by coefficients and add into the result
    this->addToInternalField(result, !add, patchId, lduAddr, coeffs, pnf);
}


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