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

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

namespace Foam
{
    makePatchFieldTypeNames(jumpCyclic);
}

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

template<>
void Foam::jumpCyclicFvPatchgpuField<Foam::scalar>::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
{
    scalargpuField 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
    (
        reinterpret_cast<const void*>(&psiInternal)
     == reinterpret_cast<const void*>(&this->internalField())
    )
    {
        gpuField<scalar> jf(this->jump());

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

        thrust::transform
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + nbrFaceCells.size(),
            jf.begin(),
            pnf.begin(),
            subtractOperatorFunctor<scalar,scalar,scalar>()
        );
    }
    else
    {
        thrust::copy
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + nbrFaceCells.size(),
            pnf.begin()
        );
    }

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

    // Multiply the field by coefficients and add into the result
    coupledFvPatchgpuField<scalar>::updateInterfaceMatrix(result, coeffs, pnf, !add);
}

template<>
void Foam::jumpCyclicFvPatchgpuField<Foam::vector>::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
{
    scalargpuField pnf(this->size());

    const labelgpuList& nbrFaceCells =
        lduAddr.patchAddr
        (
            this->cyclicPatch().neighbPatchID()
        );

    const gpuField<vector>& iField = this->internalField();

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

    // only apply jump to original field
 
    if
    (
        reinterpret_cast<const void*>(&psiInternal)
     == reinterpret_cast<const void*>(&(iField.component(cmpt).ref()))
    )
    {
        gpuField<vector> jf(this->jump());

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

        thrust::transform
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + nbrFaceCells.size(),
            jf.component(cmpt).ref().begin(),
            pnf.begin(),
            subtractOperatorFunctor<scalar,scalar,scalar>()
        );

    }
    else
    {
        thrust::copy
        (
            nbrInternalValuesStart,
            nbrInternalValuesStart + nbrFaceCells.size(),
            pnf.begin()
        );
    }

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

    // Multiply the field by coefficients and add into the result
    coupledFvPatchgpuField<vector>::updateInterfaceMatrix(result, coeffs, pnf, !add);
}

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