/* ----------------------------------------------------------------------
   LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
   http://lammps.sandia.gov, Sandia National Laboratories
   Steve Plimpton, sjplimp@sandia.gov

   Copyright (2003) Sandia Corporation.  Under the terms of Contract
   DE-AC04-94AL85000 with Sandia Corporation, the U.S. Government retains
   certain rights in this software.  This software is distributed under
   the GNU General Public License.

   See the README file in the top-level LAMMPS directory.
------------------------------------------------------------------------- */

/* ----------------------------------------------------------------------
   Contributing author: Rodrigo Canales (RWTH Aachen University)
------------------------------------------------------------------------- */

#include <math.h>
#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include "pair_buck_coul_long_intel.h"
#include "atom.h"
#include "comm.h"
#include "force.h"
#include "group.h"
#include "kspace.h"
#include "neighbor.h"
#include "neigh_list.h"
#include "neigh_request.h"
#include "math_const.h"
#include "memory.h"
#include "error.h"
#include "suffix.h"
#include "force.h"
#include "modify.h"

using namespace LAMMPS_NS;
using namespace MathConst;

#define C_FORCE_T typename ForceConst<flt_t>::c_force_t
#define C_ENERGY_T typename ForceConst<flt_t>::c_energy_t
#define TABLE_T typename ForceConst<flt_t>::table_t

PairBuckCoulLongIntel::PairBuckCoulLongIntel(LAMMPS *lmp) :
  PairBuckCoulLong(lmp)
{
  suffix_flag |= Suffix::INTEL;
}

PairBuckCoulLongIntel::~PairBuckCoulLongIntel()
{
}

void PairBuckCoulLongIntel::compute(int eflag, int vflag)
{
  if (fix->precision()==FixIntel::PREC_MODE_MIXED)
    compute<float,double>(eflag, vflag, fix->get_mixed_buffers(), 
                          force_const_single);
  else if (fix->precision()==FixIntel::PREC_MODE_DOUBLE)
    compute<double,double>(eflag, vflag, fix->get_double_buffers(),
                           force_const_double);
  else
    compute<float,float>(eflag, vflag, fix->get_single_buffers(),
                         force_const_single);

  fix->balance_stamp();
  vflag_fdotr = 0;
}

template <class flt_t, class acc_t>
void PairBuckCoulLongIntel::compute(int eflag, int vflag,
				    IntelBuffers<flt_t,acc_t> *buffers,
				    const ForceConst<flt_t> &fc)
{
  if (eflag || vflag) {
    ev_setup(eflag,vflag);
  } else evflag = vflag_fdotr = 0;

  const int inum = list->inum;
  const int nthreads = comm->nthreads;
  const int host_start = fix->host_start_pair();
  const int offload_end = fix->offload_end_pair();
  const int ago = neighbor->ago;

  if (_lrt == 0 && ago != 0 && fix->separate_buffers() == 0) {
    fix->start_watch(TIME_PACK);
    #if defined(_OPENMP)
    #pragma omp parallel default(none) shared(eflag,vflag,buffers,fc)
    #endif
    {
      int ifrom, ito, tid;
      IP_PRE_omp_range_id_align(ifrom, ito, tid, atom->nlocal + atom->nghost, 
				nthreads, sizeof(ATOM_T));
      buffers->thr_pack(ifrom,ito,ago);
    }
    fix->stop_watch(TIME_PACK);
  }
  
  if (evflag || vflag_fdotr) {
    int ovflag = 0;
    if (vflag_fdotr) ovflag = 2;
    else if (vflag) ovflag = 1;
    if (eflag) {
      if (force->newton_pair) {
	eval<1,1,1>(1, ovflag, buffers, fc, 0, offload_end);
	eval<1,1,1>(0, ovflag, buffers, fc, host_start, inum);
      } else {
	eval<1,1,0>(1, ovflag, buffers, fc, 0, offload_end);
	eval<1,1,0>(0, ovflag, buffers, fc, host_start, inum);
      }
    } else {
      if (force->newton_pair) {
	eval<1,0,1>(1, ovflag, buffers, fc, 0, offload_end);
	eval<1,0,1>(0, ovflag, buffers, fc, host_start, inum);
      } else {
	eval<1,0,0>(1, ovflag, buffers, fc, 0, offload_end);
	eval<1,0,0>(0, ovflag, buffers, fc, host_start, inum);
      }
    }
  } else {
    if (force->newton_pair) {
      eval<0,0,1>(1, 0, buffers, fc, 0, offload_end);
      eval<0,0,1>(0, 0, buffers, fc, host_start, inum);
    } else {
      eval<0,0,0>(1, 0, buffers, fc, 0, offload_end);
      eval<0,0,0>(0, 0, buffers, fc, host_start, inum);
    }
  }
}

/* ---------------------------------------------------------------------- */

template <int EVFLAG, int EFLAG, int NEWTON_PAIR, class flt_t, class acc_t>
void PairBuckCoulLongIntel::eval(const int offload, const int vflag,
				 IntelBuffers<flt_t,acc_t> *buffers,
				 const ForceConst<flt_t> &fc,
				 const int astart, const int aend)
{
  const int inum = aend - astart;
  if (inum == 0) return;
  int nlocal, nall, minlocal;
  fix->get_buffern(offload, nlocal, nall, minlocal);

  const int ago = neighbor->ago;
  IP_PRE_pack_separate_buffers(fix, buffers, ago, offload, nlocal, nall);

  ATOM_T * _noalias const x = buffers->get_x(offload);
  flt_t * _noalias const q = buffers->get_q(offload);

  const int * _noalias const numneigh = list->numneigh;
  const int * _noalias const cnumneigh = buffers->cnumneigh(list);
  const int * _noalias const firstneigh = buffers->firstneigh(list);

  const flt_t * _noalias const special_coul = fc.special_coul;
  const flt_t * _noalias const special_lj = fc.special_lj;
  const flt_t qqrd2e = force->qqrd2e;

  const C_FORCE_T * _noalias const c_force = fc.c_force[0];
  const C_ENERGY_T * _noalias const c_energy = fc.c_energy[0];
  const flt_t * _noalias const rho_inv = fc.rho_inv[0];
  const TABLE_T * _noalias const table = fc.table;
  const flt_t * _noalias const etable = fc.etable;
  const flt_t * _noalias const detable = fc.detable;
  const flt_t * _noalias const ctable = fc.ctable;
  const flt_t * _noalias const dctable = fc.dctable;
  const flt_t g_ewald = fc.g_ewald;
  const flt_t tabinnersq = fc.tabinnersq;

  const int ntypes = atom->ntypes + 1;
  const int eatom = this->eflag_atom;

  // Determine how much data to transfer
  int x_size, q_size, f_stride, ev_size, separate_flag;
  IP_PRE_get_transfern(ago, NEWTON_PAIR, EVFLAG, EFLAG, vflag,
		       buffers, offload, fix, separate_flag,
		       x_size, q_size, ev_size, f_stride);

  int tc;
  FORCE_T * _noalias f_start;
  acc_t * _noalias ev_global;
  IP_PRE_get_buffers(offload, buffers, fix, tc, f_start, ev_global);

  const int nthreads = tc;
  #ifdef _LMP_INTEL_OFFLOAD
  int *overflow = fix->get_off_overflow_flag();
  double *timer_compute = fix->off_watch_pair();
  // Redeclare as local variables for offload
  const int ncoultablebits = this->ncoultablebits;
  const int ncoulmask = this->ncoulmask;
  const int ncoulshiftbits = this->ncoulshiftbits;
  #ifdef INTEL_ALLOW_TABLE
  #define ITABLE_IN in(table,etable,detable:length(0) alloc_if(0) free_if(0)) \
                    in(ctable,dctable:length(0) alloc_if(0) free_if(0)) \
                    in(ncoultablebits,tabinnersq,ncoulmask,ncoulshiftbits)
  #else
  #define ITABLE_IN
  #endif

  if (offload) fix->start_watch(TIME_OFFLOAD_LATENCY);
  #pragma offload target(mic:_cop) if(offload)                 \
    in(special_lj,special_coul:length(0) alloc_if(0) free_if(0)) \
    in(c_force, c_energy:length(0) alloc_if(0) free_if(0)) \
    in(rho_inv:length(0) alloc_if(0) free_if(0)) \
    in(firstneigh:length(0) alloc_if(0) free_if(0)) \
    in(cnumneigh:length(0) alloc_if(0) free_if(0)) \
    in(numneigh:length(0) alloc_if(0) free_if(0)) \
    in(x:length(x_size) alloc_if(0) free_if(0)) \
    in(q:length(q_size) alloc_if(0) free_if(0)) \
    in(overflow:length(0) alloc_if(0) free_if(0)) \
    in(astart,nthreads,qqrd2e,g_ewald,inum,nall,ntypes,vflag,eatom) \
    in(f_stride,nlocal,minlocal,separate_flag,offload) \
    out(f_start:length(f_stride) alloc_if(0) free_if(0)) \
    out(ev_global:length(ev_size) alloc_if(0) free_if(0)) \
    out(timer_compute:length(1) alloc_if(0) free_if(0)) \
    ITABLE_IN signal(f_start)
  #endif
  {
    #if defined(__MIC__) && defined(_LMP_INTEL_OFFLOAD)
    *timer_compute = MIC_Wtime();
    #endif

    IP_PRE_repack_for_offload(NEWTON_PAIR, separate_flag, nlocal, nall, 
			      f_stride, x, q);

    acc_t oevdwl, oecoul, ov0, ov1, ov2, ov3, ov4, ov5;
    if (EVFLAG) {
      oevdwl = oecoul = (acc_t)0;
      if (vflag) ov0 = ov1 = ov2 = ov3 = ov4 = ov5 = (acc_t)0;
    }

    // loop over neighbors of my atoms
    #if defined(_OPENMP)
    #pragma omp parallel default(none)        \
      shared(f_start,f_stride,nlocal,nall,minlocal)	\
      reduction(+:oevdwl,oecoul,ov0,ov1,ov2,ov3,ov4,ov5)
    #endif
    {
      int iifrom, iito, tid;
      IP_PRE_omp_range_id(iifrom, iito, tid, inum, nthreads);
      iifrom += astart;
      iito += astart;

      FORCE_T * _noalias const f = f_start - minlocal + (tid * f_stride);
      memset(f + minlocal, 0, f_stride * sizeof(FORCE_T));

      for (int i = iifrom; i < iito; ++i) {
        const int itype = x[i].w;
        const int ptr_off = itype * ntypes;
        const C_FORCE_T * _noalias const c_forcei = c_force + ptr_off;
        const C_ENERGY_T * _noalias const c_energyi = c_energy + ptr_off;
        const flt_t * _noalias const rho_invi = rho_inv + ptr_off; 

        const int   * _noalias const jlist = firstneigh + cnumneigh[i];
        const int jnum = numneigh[i];

        acc_t fxtmp,fytmp,fztmp,fwtmp;
	acc_t sevdwl, secoul, sv0, sv1, sv2, sv3, sv4, sv5;

        const flt_t xtmp = x[i].x;
        const flt_t ytmp = x[i].y;
        const flt_t ztmp = x[i].z;
        const flt_t qtmp = q[i];
        fxtmp = fytmp = fztmp = (acc_t)0;
        if (EVFLAG) {
	  if (EFLAG) fwtmp = sevdwl = secoul = (acc_t)0;
	  if (vflag==1) sv0 = sv1 = sv2 = sv3 = sv4 = sv5 = (acc_t)0;
	}

        #if defined(LMP_SIMD_COMPILER)
        #pragma vector aligned
	#pragma simd reduction(+:fxtmp, fytmp, fztmp, fwtmp, sevdwl, \
	                       sv0, sv1, sv2, sv3, sv4, sv5)
        #endif
        for (int jj = 0; jj < jnum; jj++) {
          flt_t forcecoul, forcebuck, evdwl, ecoul;
          forcecoul = forcebuck = evdwl = ecoul = (flt_t)0.0;

          const int sbindex = jlist[jj] >> SBBITS & 3;
          const int j = jlist[jj] & NEIGHMASK;

          const flt_t delx = xtmp - x[j].x;
          const flt_t dely = ytmp - x[j].y;
          const flt_t delz = ztmp - x[j].z;
          const int jtype = x[j].w;
          const flt_t rsq = delx * delx + dely * dely + delz * delz;
          const flt_t r2inv = (flt_t)1.0 / rsq;
          const flt_t r = (flt_t)1.0 / sqrt(r2inv);

          #ifdef INTEL_VMASK
          if (rsq < c_forcei[jtype].cutsq) {
          #endif
            #ifdef INTEL_ALLOW_TABLE
            if (!ncoultablebits || rsq <= tabinnersq) {
            #endif
              const flt_t A1 =  0.254829592;
              const flt_t A2 = -0.284496736;
              const flt_t A3 =  1.421413741;
              const flt_t A4 = -1.453152027;
              const flt_t A5 =  1.061405429;
              const flt_t EWALD_F = 1.12837917;
              const flt_t INV_EWALD_P = 1.0 / 0.3275911;

              const flt_t grij = g_ewald * r;
              const flt_t expm2 = exp(-grij * grij);
              const flt_t t = INV_EWALD_P / (INV_EWALD_P + grij);
              const flt_t erfc = t * (A1+t*(A2+t*(A3+t*(A4+t*A5)))) * expm2;
              const flt_t prefactor = qqrd2e * qtmp * q[j] / r;
              forcecoul = prefactor * (erfc + EWALD_F * grij * expm2);
              if (EFLAG) ecoul = prefactor * erfc;

	      const flt_t adjust = ((flt_t)1.0 - special_coul[sbindex])*
		prefactor;
	      forcecoul -= adjust;
	      if (EFLAG) ecoul -= adjust;

            #ifdef INTEL_ALLOW_TABLE
            } else {
              float rsq_lookup = rsq;
              const int itable = (__intel_castf32_u32(rsq_lookup) &
                                  ncoulmask) >> ncoulshiftbits;
              const flt_t fraction = (rsq_lookup - table[itable].r) *
                table[itable].dr;

              const flt_t tablet = table[itable].f +
                fraction * table[itable].df;
              forcecoul = qtmp * q[j] * tablet;
              if (EFLAG) ecoul = qtmp * q[j] * (etable[itable] +
                                                fraction * detable[itable]);
              if (sbindex) {
                const flt_t table2 = ctable[itable] +
                  fraction * dctable[itable];
                const flt_t prefactor = qtmp * q[j] * table2;
                const flt_t adjust = ((flt_t)1.0 - special_coul[sbindex]) *
                  prefactor;
                forcecoul -= adjust;
                if (EFLAG) ecoul -= adjust;
              }
            }
            #endif
            #ifdef INTEL_VMASK
          }
	  #endif

	  #ifdef INTEL_VMASK
          if (rsq < c_forcei[jtype].cut_ljsq) {
          #endif
            flt_t r6inv = r2inv * r2inv * r2inv;
            flt_t rexp = exp(-r * rho_invi[jtype]);
            forcebuck = r * rexp * c_forcei[jtype].buck1 -
              r6inv * c_forcei[jtype].buck2;
            if (EFLAG) evdwl = rexp * c_energyi[jtype].a -
                         r6inv * c_energyi[jtype].c -
                         c_energyi[jtype].offset;

            if (sbindex) {
              const flt_t factor_lj = special_lj[sbindex];
              forcebuck *= factor_lj;
              if (EFLAG) evdwl *= factor_lj;
            }
          #ifdef INTEL_VMASK
          }
          #else
          if (rsq > c_forcei[jtype].cutsq)
            { forcecoul = (flt_t)0.0; ecoul = (flt_t)0.0; }
          if (rsq > c_forcei[jtype].cut_ljsq)
            { forcebuck = (flt_t)0.0; evdwl = (flt_t)0.0; }
          #endif

          #ifdef INTEL_VMASK
          if (rsq < c_forcei[jtype].cutsq) {
          #endif
            const flt_t fpair = (forcecoul + forcebuck) * r2inv;
            fxtmp += delx * fpair;
            fytmp += dely * fpair;
            fztmp += delz * fpair;
            if (NEWTON_PAIR || j < nlocal) {
              f[j].x -= delx * fpair;
              f[j].y -= dely * fpair;
              f[j].z -= delz * fpair;
            }

            if (EVFLAG) {
              flt_t ev_pre = (flt_t)0;
              if (NEWTON_PAIR || i < nlocal)
                ev_pre += (flt_t)0.5;
              if (NEWTON_PAIR || j < nlocal)
                ev_pre += (flt_t)0.5;

              if (EFLAG) {
                sevdwl += ev_pre * evdwl;
                secoul += ev_pre * ecoul;
                if (eatom) {
                  if (NEWTON_PAIR || i < nlocal)
                    fwtmp += (flt_t)0.5 * evdwl + (flt_t)0.5 * ecoul;
                  if (NEWTON_PAIR || j < nlocal) 
                    f[j].w += (flt_t)0.5 * evdwl + (flt_t)0.5 * ecoul;
                }
              }
              IP_PRE_ev_tally_nbor(vflag, ev_pre, fpair, delx, dely, delz);
            }
          #ifdef INTEL_VMASK
          }
          #endif
        } // for jj

        f[i].x += fxtmp;
        f[i].y += fytmp;
        f[i].z += fztmp;
	IP_PRE_ev_tally_atomq(EVFLAG, EFLAG, vflag, f, fwtmp);
      } // for ii

      #ifndef _LMP_INTEL_OFFLOAD
      if (vflag == 2)
      #endif
      {
        #if defined(_OPENMP)
        #pragma omp barrier
        #endif
        IP_PRE_fdotr_acc_force(NEWTON_PAIR, EVFLAG,  EFLAG, vflag, eatom, nall,
			       nlocal, minlocal, nthreads, f_start, f_stride, 
			       x, offload);
      }
    } // end of omp parallel region
    if (EVFLAG) {
      if (EFLAG) {
        ev_global[0] = oevdwl;
        ev_global[1] = oecoul;
      }
      if (vflag) {
        ev_global[2] = ov0;
        ev_global[3] = ov1;
        ev_global[4] = ov2;
        ev_global[5] = ov3;
        ev_global[6] = ov4;
        ev_global[7] = ov5;
      }
    }
    #if defined(__MIC__) && defined(_LMP_INTEL_OFFLOAD)
    *timer_compute = MIC_Wtime() - *timer_compute;
    #endif
  } // end of offload region

  if (offload)
    fix->stop_watch(TIME_OFFLOAD_LATENCY);
  else
    fix->stop_watch(TIME_HOST_PAIR);

  if (EVFLAG)
    fix->add_result_array(f_start, ev_global, offload, eatom, 0, vflag);
  else
    fix->add_result_array(f_start, 0, offload);
}

/* ---------------------------------------------------------------------- */

void PairBuckCoulLongIntel::init_style()
{
  PairBuckCoulLong::init_style();
  neighbor->requests[neighbor->nrequest-1]->intel = 1;

  int ifix = modify->find_fix("package_intel");
  if (ifix < 0)
    error->all(FLERR,
               "The 'package intel' command is required for /intel styles");
  fix = static_cast<FixIntel *>(modify->fix[ifix]);
  
  fix->pair_init_check();
  #ifdef _LMP_INTEL_OFFLOAD
  _cop = fix->coprocessor_number();
  #endif

  if (fix->precision() == FixIntel::PREC_MODE_MIXED)
    pack_force_const(force_const_single, fix->get_mixed_buffers());
  else if (fix->precision() == FixIntel::PREC_MODE_DOUBLE)
    pack_force_const(force_const_double, fix->get_double_buffers());
  else
    pack_force_const(force_const_single, fix->get_single_buffers());

  _lrt = fix->lrt();
}

template <class flt_t, class acc_t>
void PairBuckCoulLongIntel::pack_force_const(ForceConst<flt_t> &fc,
                                          IntelBuffers<flt_t,acc_t> *buffers)
{
  int tp1 = atom->ntypes + 1;
  int ntable = 1;
  if (ncoultablebits)
    for (int i = 0; i < ncoultablebits; i++) ntable *= 2;

  fc.set_ntypes(tp1, ntable, memory, _cop);
  buffers->set_ntypes(tp1);
  flt_t **cutneighsq = buffers->get_cutneighsq();

  // Repeat cutsq calculation because done after call to init_style
  double cut, cutneigh;
  for (int i = 1; i <= atom->ntypes; i++) {
    for (int j = i; j <= atom->ntypes; j++) {
      if (setflag[i][j] != 0 || (setflag[i][i] != 0 && setflag[j][j] != 0)) {
        cut = init_one(i, j);
        cutneigh = cut + neighbor->skin;
        cutsq[i][j] = cutsq[j][i] = cut*cut;
        cutneighsq[i][j] = cutneighsq[j][i] = cutneigh * cutneigh;
      }
    }
  }

  fc.g_ewald = force->kspace->g_ewald;
  fc.tabinnersq = tabinnersq;

  for (int i = 0; i < 4; i++) {
    fc.special_lj[i] = force->special_lj[i];
    fc.special_coul[i] = force->special_coul[i];
    fc.special_coul[0] = 1.0;
    fc.special_lj[0] = 1.0;
  }

  for (int i = 0; i < tp1; i++) {
    for (int j = 0; j < tp1; j++) {
      fc.c_force[i][j].cutsq = cutsq[i][j];
      fc.c_force[i][j].cut_ljsq = cut_ljsq[i][j];
      fc.c_force[i][j].buck1 = buck1[i][j];
      fc.c_force[i][j].buck2 = buck2[i][j];
      fc.rho_inv[i][j] = rhoinv[i][j];
      fc.c_energy[i][j].a = a[i][j];
      fc.c_energy[i][j].c = c[i][j];
      fc.c_energy[i][j].offset = offset[i][j];
      fc.c_energy[i][j].pad = rhoinv[i][j];
    }
  }

  if (ncoultablebits) {
    for (int i = 0; i < ntable; i++) {
      fc.table[i].r = rtable[i];
      fc.table[i].dr = drtable[i];
      fc.table[i].f = ftable[i];
      fc.table[i].df = dftable[i];
      fc.etable[i] = etable[i];
      fc.detable[i] = detable[i];
      fc.ctable[i] = ctable[i];
      fc.dctable[i] = dctable[i];
    }
  }

  #ifdef _LMP_INTEL_OFFLOAD
  if (_cop < 0) return;
  flt_t * special_lj = fc.special_lj;
  flt_t * special_coul = fc.special_coul;
  C_FORCE_T * c_force = fc.c_force[0];
  C_ENERGY_T * c_energy = fc.c_energy[0];
  TABLE_T * table = fc.table;
  flt_t * rho_inv = fc.rho_inv[0];
  flt_t * etable = fc.etable;
  flt_t * detable = fc.detable;
  flt_t * ctable = fc.ctable;
  flt_t * dctable = fc.dctable;
  flt_t * ocutneighsq = cutneighsq[0];
  int tp1sq = tp1 * tp1;
  #pragma offload_transfer target(mic:_cop) \
    in(special_lj, special_coul: length(4) alloc_if(0) free_if(0)) \
    in(c_force, c_energy: length(tp1sq) alloc_if(0) free_if(0)) \
    in(rho_inv: length(tp1sq) alloc_if(0) free_if(0)) \
    in(table: length(ntable) alloc_if(0) free_if(0)) \
    in(etable,detable,ctable,dctable: length(ntable) alloc_if(0) free_if(0)) \
    in(ocutneighsq: length(tp1sq) alloc_if(0) free_if(0))
  #endif
}

/* ---------------------------------------------------------------------- */

template <class flt_t>
void PairBuckCoulLongIntel::ForceConst<flt_t>::set_ntypes(const int ntypes,
							   const int ntable,
							   Memory *memory,
							   const int cop) {
  if ( (ntypes != _ntypes || ntable != _ntable) ) {
    if (_ntypes > 0) {
      #ifdef _LMP_INTEL_OFFLOAD
      flt_t * ospecial_lj = special_lj;
      flt_t * ospecial_coul = special_coul;
      c_force_t * oc_force = c_force[0];
      c_energy_t * oc_energy = c_energy[0];
      table_t * otable = table;
      flt_t * orho_inv = rho_inv[0];
      flt_t * oetable = etable;
      flt_t * odetable = detable;
      flt_t * octable = ctable;
      flt_t * odctable = dctable;
      if (ospecial_lj != NULL && oc_force != NULL && orho_inv != NULL &&
          oc_energy != NULL && otable != NULL && oetable != NULL &&
          odetable != NULL && octable != NULL && odctable != NULL &&
          ospecial_coul != NULL && _cop >= 0) {
        #pragma offload_transfer target(mic:cop) \
          nocopy(ospecial_lj, ospecial_coul: alloc_if(0) free_if(1)) \
	  nocopy(oc_force, oc_energy: alloc_if(0) free_if(1)) \
	  nocopy(orho_inv: alloc_if(0) free_if(1)) \
	  nocopy(otable: alloc_if(0) free_if(1)) \
	  nocopy(oetable, odetable, octable, odctable: alloc_if(0) free_if(1))
      }
      #endif

      _memory->destroy(c_force);
      _memory->destroy(c_energy);
      _memory->destroy(table);
      _memory->destroy(rho_inv);
      _memory->destroy(etable);
      _memory->destroy(detable);
      _memory->destroy(ctable);
      _memory->destroy(dctable);
    }
    if (ntypes > 0) {
      _cop = cop;
      memory->create(c_force,ntypes,ntypes,"fc.c_force");
      memory->create(c_energy,ntypes,ntypes,"fc.c_energy");
      memory->create(rho_inv,ntypes,ntypes,"fc.rho_inv");
      memory->create(table,ntable,"pair:fc.table");
      memory->create(etable,ntable,"pair:fc.etable");
      memory->create(detable,ntable,"pair:fc.detable");
      memory->create(ctable,ntable,"pair:fc.ctable");
      memory->create(dctable,ntable,"pair:fc.dctable");

      #ifdef _LMP_INTEL_OFFLOAD
      flt_t * ospecial_lj = special_lj;
      flt_t * ospecial_coul = special_coul;
      c_force_t * oc_force = c_force[0];
      c_energy_t * oc_energy = c_energy[0];
      table_t * otable = table;
      flt_t * orho_inv = rho_inv[0];
      flt_t * oetable = etable;
      flt_t * odetable = detable;
      flt_t * octable = ctable;
      flt_t * odctable = dctable;
      int tp1sq = ntypes*ntypes;
      if (ospecial_lj != NULL && oc_force != NULL && orho_inv != NULL &&
          oc_energy != NULL && otable !=NULL && oetable != NULL &&
          odetable != NULL && octable != NULL && odctable != NULL &&
          ospecial_coul != NULL && cop >= 0) {
        #pragma offload_transfer target(mic:cop) \
          nocopy(ospecial_lj: length(4) alloc_if(1) free_if(0)) \
          nocopy(ospecial_coul: length(4) alloc_if(1) free_if(0)) \
          nocopy(oc_force: length(tp1sq) alloc_if(1) free_if(0)) \
          nocopy(oc_energy: length(tp1sq) alloc_if(1) free_if(0)) \
          nocopy(orho_inv: length(tp1sq) alloc_if(1) free_if(0)) \
          nocopy(otable: length(ntable) alloc_if(1) free_if(0)) \
          nocopy(oetable,odetable: length(ntable) alloc_if(1) free_if(0)) \
          nocopy(octable,odctable: length(ntable) alloc_if(1) free_if(0))
      }
      #endif
    }
  }
  _ntypes=ntypes;
  _ntable=ntable;
  _memory=memory;
}


