diff --git a/lib/cuda/pair_lj_charmm_coul_charmm_cuda_kernel_nc.cu b/lib/cuda/pair_lj_charmm_coul_charmm_cuda_kernel_nc.cu
index d741e7ac9..cc1963bfc 100644
--- a/lib/cuda/pair_lj_charmm_coul_charmm_cuda_kernel_nc.cu
+++ b/lib/cuda/pair_lj_charmm_coul_charmm_cuda_kernel_nc.cu
@@ -1,73 +1,71 @@
 /* ----------------------------------------------------------------------
    LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
 
    Original Version:
    http://lammps.sandia.gov, Sandia National Laboratories
    Steve Plimpton, sjplimp@sandia.gov
 
    See the README file in the top-level LAMMPS directory.
 
    -----------------------------------------------------------------------
 
    USER-CUDA Package and associated modifications:
    https://sourceforge.net/projects/lammpscuda/
 
    Christian Trott, christian.trott@tu-ilmenau.de
    Lars Winterfeld, lars.winterfeld@tu-ilmenau.de
    Theoretical Physics II, University of Technology Ilmenau, Germany
 
    See the README file in the USER-CUDA directory.
 
    This software is distributed under the GNU General Public License.
 ------------------------------------------------------------------------- */
 __device__ inline F_CFLOAT PairLJCharmmCuda_Eval(const F_CFLOAT &rsq, const int ij_type, F_CFLOAT &factor_lj, int &eflag, ENERGY_CFLOAT &evdwl)
 {
   const F_CFLOAT r2inv = F_F(1.0) / rsq;
   const F_CFLOAT r6inv = r2inv * r2inv * r2inv;
   F_CFLOAT forcelj = r6inv * (_lj1[ij_type] * r6inv - _lj2[ij_type]);
   F_CFLOAT philj, switch1;
 
   if(rsq > _cut_innersq_global) {
     switch1 = (_cutsq_global - rsq) * (_cutsq_global - rsq) *
               (_cutsq_global + F_F(2.0) * rsq - F_F(3.0) * _cut_innersq_global) * _denom_lj_inv;
     const F_CFLOAT switch2 = F_F(12.0) * rsq * (_cutsq_global - rsq) *
                             (rsq - _cut_innersq_global) * _denom_lj_inv;
     philj = r6inv * (_lj3[ij_type] * r6inv - _lj4[ij_type]);
     forcelj = forcelj * switch1 + philj * switch2;
   }
 
   if(eflag) {
     ENERGY_CFLOAT evdwl_tmp = factor_lj;
 
     if(rsq > _cut_innersq_global) {
       evdwl_tmp *= philj * switch1;
     } else
       evdwl_tmp *= r6inv * (_lj3[ij_type] * r6inv - _lj4[ij_type]);
 
     evdwl += evdwl_tmp;
   }
 
   return factor_lj * forcelj * r2inv;
 }
 
 __device__ inline F_CFLOAT CoulCharmmCuda_Eval(const F_CFLOAT &rsq, F_CFLOAT &factor_coul, int &eflag, ENERGY_CFLOAT &ecoul, F_CFLOAT qij)
 {
   F_CFLOAT forcecoul;
   ENERGY_CFLOAT ecoul_tmp = forcecoul = _qqrd2e * qij * _RSQRT_(rsq) * factor_coul;
 
   if(rsq > _cut_coul_innersq_global) {
     const F_CFLOAT switch1 = (_cut_coulsq_global - rsq) * (_cut_coulsq_global - rsq) *
                             (_cut_coulsq_global + F_F(2.0) * rsq - F_F(3.0) * _cut_coul_innersq_global) * _denom_coul_inv;
     ecoul_tmp *= switch1;
-    const F_CFLOAT switch2 = F_F(12.0) * rsq * (_cut_coulsq_global - rsq) *
-                            (rsq - _cut_coul_innersq_global) * _denom_coul_inv;
-    forcecoul *= switch1 + switch2;
+    forcecoul *= switch1;
   }
 
   if(eflag) {
     ecoul += ecoul_tmp * factor_coul;
   }
 
   return forcecoul * (F_F(1.0) / rsq);
 }
 
diff --git a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp
index 98f9e8237..cd1d7a02d 100644
--- a/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp
+++ b/src/KOKKOS/pair_lj_charmm_coul_charmm_kokkos.cpp
@@ -1,519 +1,518 @@
 /* ----------------------------------------------------------------------
    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: Ray Shan (SNL)
 ------------------------------------------------------------------------- */
 
 #include <math.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
 #include "pair_lj_charmm_coul_charmm_kokkos.h"
 #include "kokkos.h"
 #include "atom_kokkos.h"
 #include "comm.h"
 #include "force.h"
 #include "neighbor.h"
 #include "neigh_list.h"
 #include "neigh_request.h"
 #include "update.h"
 #include "integrate.h"
 #include "respa.h"
 #include "math_const.h"
 #include "memory.h"
 #include "error.h"
 #include "atom_masks.h"
 
 using namespace LAMMPS_NS;
 using namespace MathConst;
 
 #define KOKKOS_CUDA_MAX_THREADS 256
 #define KOKKOS_CUDA_MIN_BLOCKS 8
 
 
 #define EWALD_F   1.12837917
 #define EWALD_P   0.3275911
 #define A1        0.254829592
 #define A2       -0.284496736
 #define A3        1.421413741
 #define A4       -1.453152027
 #define A5        1.061405429
 
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType>
 PairLJCharmmCoulCharmmKokkos<DeviceType>::PairLJCharmmCoulCharmmKokkos(LAMMPS *lmp):PairLJCharmmCoulCharmm(lmp)
 {
   respa_enable = 0;
 
   atomKK = (AtomKokkos *) atom;
   execution_space = ExecutionSpaceFromDevice<DeviceType>::space;
   datamask_read = X_MASK | F_MASK | TYPE_MASK | Q_MASK | ENERGY_MASK | VIRIAL_MASK;
   datamask_modify = F_MASK | ENERGY_MASK | VIRIAL_MASK;
   cutsq = NULL;
   cut_ljsq = 0.0;
   cut_coulsq = 0.0;
 
 }
 
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType>
 PairLJCharmmCoulCharmmKokkos<DeviceType>::~PairLJCharmmCoulCharmmKokkos()
 {
   if (!copymode) {
     memory->destroy_kokkos(k_eatom,eatom);
     memory->destroy_kokkos(k_vatom,vatom);
     k_cutsq = DAT::tdual_ffloat_2d();
     k_cut_ljsq = DAT::tdual_ffloat_2d();
     k_cut_coulsq = DAT::tdual_ffloat_2d();
     memory->sfree(cutsq);
     //memory->sfree(cut_ljsq);
     //memory->sfree(cut_coulsq);
     eatom = NULL;
     vatom = NULL;
     cutsq = NULL;
     cut_ljsq = 0.0;
     cut_coulsq = 0.0;
   }
 }
 
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::cleanup_copy() {
   // WHY needed: this prevents parent copy from deallocating any arrays
   allocated = 0;
   cutsq = NULL;
   cut_ljsq = 0.0;
   eatom = NULL;
   vatom = NULL;
   ftable = NULL;
 }
 
 /* ---------------------------------------------------------------------- */
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::compute(int eflag_in, int vflag_in)
 {
   eflag = eflag_in;
   vflag = vflag_in;
 
   if (neighflag == FULL || neighflag == FULLCLUSTER) no_virial_fdotr_compute = 1;
 
   if (eflag || vflag) ev_setup(eflag,vflag);
   else evflag = vflag_fdotr = 0;
 
   atomKK->sync(execution_space,datamask_read);
   k_cutsq.template sync<DeviceType>();
   k_cut_ljsq.template sync<DeviceType>();
   k_cut_coulsq.template sync<DeviceType>();
   k_params.template sync<DeviceType>();
   if (eflag || vflag) atomKK->modified(execution_space,datamask_modify);
   else atomKK->modified(execution_space,F_MASK);
 
   x = atomKK->k_x.view<DeviceType>();
   c_x = atomKK->k_x.view<DeviceType>();
   f = atomKK->k_f.view<DeviceType>();
   q = atomKK->k_q.view<DeviceType>();
   type = atomKK->k_type.view<DeviceType>();
   nlocal = atom->nlocal;
   nall = atom->nlocal + atom->nghost;
   special_lj[0] = force->special_lj[0];
   special_lj[1] = force->special_lj[1];
   special_lj[2] = force->special_lj[2];
   special_lj[3] = force->special_lj[3];
   special_coul[0] = force->special_coul[0];
   special_coul[1] = force->special_coul[1];
   special_coul[2] = force->special_coul[2];
   special_coul[3] = force->special_coul[3];
   qqrd2e = force->qqrd2e;
   newton_pair = force->newton_pair;
 
   // loop over neighbors of my atoms
 
   copymode = 1;
 
   EV_FLOAT ev;
   if(ncoultablebits)
     ev = pair_compute<PairLJCharmmCoulCharmmKokkos<DeviceType>,CoulLongTable<1> >
       (this,(NeighListKokkos<DeviceType>*)list);
   else
     ev = pair_compute<PairLJCharmmCoulCharmmKokkos<DeviceType>,CoulLongTable<0> >
       (this,(NeighListKokkos<DeviceType>*)list);
 
 
   if (eflag) {
     eng_vdwl += ev.evdwl;
     eng_coul += ev.ecoul;
   }
   if (vflag_global) {
     virial[0] += ev.v[0];
     virial[1] += ev.v[1];
     virial[2] += ev.v[2];
     virial[3] += ev.v[3];
     virial[4] += ev.v[4];
     virial[5] += ev.v[5];
   }
 
   if (vflag_fdotr) pair_virial_fdotr_compute(this);
 
   copymode = 0;
 }
 
 /* ----------------------------------------------------------------------
    compute LJ CHARMM pair force between atoms i and j
    ---------------------------------------------------------------------- */
 template<class DeviceType>
 template<bool STACKPARAMS, class Specialisation>
 KOKKOS_INLINE_FUNCTION
 F_FLOAT PairLJCharmmCoulCharmmKokkos<DeviceType>::
 compute_fpair(const F_FLOAT& rsq, const int& i, const int&j,
               const int& itype, const int& jtype) const {
   const F_FLOAT r2inv = 1.0/rsq;
   const F_FLOAT r6inv = r2inv*r2inv*r2inv;
   F_FLOAT forcelj, switch1, switch2, englj;
 
   forcelj = r6inv *
     ((STACKPARAMS?m_params[itype][jtype].lj1:params(itype,jtype).lj1)*r6inv -
      (STACKPARAMS?m_params[itype][jtype].lj2:params(itype,jtype).lj2));
 
   if (rsq > cut_lj_innersq) {
     switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
               (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
     switch2 = 12.0*rsq * (cut_ljsq-rsq) * (rsq-cut_lj_innersq) / denom_lj;
     englj = r6inv *
 	    ((STACKPARAMS?m_params[itype][jtype].lj3:params(itype,jtype).lj3)*r6inv -
 	     (STACKPARAMS?m_params[itype][jtype].lj4:params(itype,jtype).lj4));
     forcelj = forcelj*switch1 + englj*switch2;
   }
 
   return forcelj*r2inv;
 }
 
 /* ----------------------------------------------------------------------
    compute LJ CHARMM pair potential energy between atoms i and j
    ---------------------------------------------------------------------- */
 template<class DeviceType>
 template<bool STACKPARAMS, class Specialisation>
 KOKKOS_INLINE_FUNCTION
 F_FLOAT PairLJCharmmCoulCharmmKokkos<DeviceType>::
 compute_evdwl(const F_FLOAT& rsq, const int& i, const int&j,
               const int& itype, const int& jtype) const {
   const F_FLOAT r2inv = 1.0/rsq;
   const F_FLOAT r6inv = r2inv*r2inv*r2inv;
   F_FLOAT englj, switch1;
 
   englj = r6inv *
     ((STACKPARAMS?m_params[itype][jtype].lj3:params(itype,jtype).lj3)*r6inv -
      (STACKPARAMS?m_params[itype][jtype].lj4:params(itype,jtype).lj4));
 
   if (rsq > cut_lj_innersq) {
     switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
       (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
     englj *= switch1;
   }
 
   return englj;
 
 }
 
 /* ----------------------------------------------------------------------
    compute coulomb pair force between atoms i and j
    ---------------------------------------------------------------------- */
 template<class DeviceType>
 template<bool STACKPARAMS,  class Specialisation>
 KOKKOS_INLINE_FUNCTION
 F_FLOAT PairLJCharmmCoulCharmmKokkos<DeviceType>::
 compute_fcoul(const F_FLOAT& rsq, const int& i, const int&j,
               const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const {
 
   const F_FLOAT r2inv = 1.0/rsq;
   const F_FLOAT rinv = sqrt(r2inv);
-  F_FLOAT forcecoul, switch1, switch2;
+  F_FLOAT forcecoul, switch1;
 
   forcecoul = qqrd2e*qtmp*q(j) *rinv;
 
   if (rsq > cut_coul_innersq) {
     switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
               (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) / denom_coul;
-    switch2 = 12.0*rsq * (cut_coulsq-rsq) * (rsq-cut_coul_innersq) / denom_coul;
-    forcecoul *= switch1 + switch2;
+    forcecoul *= switch1;
   }
 
   return forcecoul * r2inv * factor_coul;
 
 }
 
 /* ----------------------------------------------------------------------
    compute coulomb pair potential energy between atoms i and j
    ---------------------------------------------------------------------- */
 template<class DeviceType>
 template<bool STACKPARAMS, class Specialisation>
 KOKKOS_INLINE_FUNCTION
 F_FLOAT PairLJCharmmCoulCharmmKokkos<DeviceType>::
 compute_ecoul(const F_FLOAT& rsq, const int& i, const int&j,
               const int& itype, const int& jtype, const F_FLOAT& factor_coul, const F_FLOAT& qtmp) const {
 
   const F_FLOAT r2inv = 1.0/rsq;
   const F_FLOAT rinv = sqrt(r2inv);
   F_FLOAT ecoul, switch1;
 
   ecoul = qqrd2e * qtmp * q(j) * rinv;
   if (rsq > cut_coul_innersq) {
     switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
               (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) /
               denom_coul;
     ecoul *= switch1;
   }
 
   return ecoul * factor_coul;
 
 }
 
 /* ----------------------------------------------------------------------
    allocate all arrays
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::allocate()
 {
   PairLJCharmmCoulCharmm::allocate();
 
   int n = atom->ntypes;
 
   memory->destroy(cutsq);
   memory->create_kokkos(k_cutsq,cutsq,n+1,n+1,"pair:cutsq");
   d_cutsq = k_cutsq.template view<DeviceType>();
 
   //memory->destroy(cut_ljsq);
   memory->create_kokkos(k_cut_ljsq,n+1,n+1,"pair:cut_ljsq");
   d_cut_ljsq = k_cut_ljsq.template view<DeviceType>();
 
   memory->create_kokkos(k_cut_coulsq,n+1,n+1,"pair:cut_coulsq");
   d_cut_coulsq = k_cut_coulsq.template view<DeviceType>();
 
   k_params = Kokkos::DualView<params_lj_coul**,Kokkos::LayoutRight,DeviceType>("PairLJCharmmCoulCharmm::params",n+1,n+1);
   params = k_params.d_view;
 }
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::init_tables(double cut_coul, double *cut_respa)
 {
   Pair::init_tables(cut_coul,cut_respa);
 
   typedef typename ArrayTypes<DeviceType>::t_ffloat_1d table_type;
   typedef typename ArrayTypes<LMPHostType>::t_ffloat_1d host_table_type;
 
   int ntable = 1;
   for (int i = 0; i < ncoultablebits; i++) ntable *= 2;
 
 
   // Copy rtable and drtable
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
   for(int i = 0; i < ntable; i++) {
     h_table(i) = rtable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_rtable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
   for(int i = 0; i < ntable; i++) {
     h_table(i) = drtable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_drtable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   // Copy ftable and dftable
   for(int i = 0; i < ntable; i++) {
     h_table(i) = ftable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_ftable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   for(int i = 0; i < ntable; i++) {
     h_table(i) = dftable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_dftable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   // Copy ctable and dctable
   for(int i = 0; i < ntable; i++) {
     h_table(i) = ctable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_ctable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   for(int i = 0; i < ntable; i++) {
     h_table(i) = dctable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_dctable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   // Copy etable and detable
   for(int i = 0; i < ntable; i++) {
     h_table(i) = etable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_etable = d_table;
   }
 
   {
   host_table_type h_table("HostTable",ntable);
   table_type d_table("DeviceTable",ntable);
 
   for(int i = 0; i < ntable; i++) {
     h_table(i) = detable[i];
   }
   Kokkos::deep_copy(d_table,h_table);
   d_detable = d_table;
   }
 }
 
 
 /* ----------------------------------------------------------------------
    global settings
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::settings(int narg, char **arg)
 {
   if (narg > 2) error->all(FLERR,"Illegal pair_style command");
 
   PairLJCharmmCoulCharmm::settings(narg,arg);
 }
 
 /* ----------------------------------------------------------------------
    init specific to this pair style
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
 void PairLJCharmmCoulCharmmKokkos<DeviceType>::init_style()
 {
   PairLJCharmmCoulCharmm::init_style();
 
   // error if rRESPA with inner levels
 
   if (update->whichflag == 1 && strstr(update->integrate_style,"respa")) {
     int respa = 0;
     if (((Respa *) update->integrate)->level_inner >= 0) respa = 1;
     if (((Respa *) update->integrate)->level_middle >= 0) respa = 2;
     if (respa)
       error->all(FLERR,"Cannot use Kokkos pair style with rRESPA inner/middle");
   }
 
   // irequest = neigh request made by parent class
 
   neighflag = lmp->kokkos->neighflag;
   int irequest = neighbor->nrequest - 1;
 
   neighbor->requests[irequest]->
     kokkos_host = Kokkos::Impl::is_same<DeviceType,LMPHostType>::value &&
     !Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
   neighbor->requests[irequest]->
     kokkos_device = Kokkos::Impl::is_same<DeviceType,LMPDeviceType>::value;
 
   if (neighflag == FULL) {
     neighbor->requests[irequest]->full = 1;
     neighbor->requests[irequest]->half = 0;
     neighbor->requests[irequest]->full_cluster = 0;
   } else if (neighflag == HALF || neighflag == HALFTHREAD) {
     neighbor->requests[irequest]->full = 0;
     neighbor->requests[irequest]->half = 1;
     neighbor->requests[irequest]->full_cluster = 0;
   } else {
     error->all(FLERR,"Cannot use chosen neighbor list style with lj/charmm/coul/charmm/kk");
   }
 }
 
 /* ----------------------------------------------------------------------
    init for one type pair i,j and corresponding j,i
 ------------------------------------------------------------------------- */
 
 template<class DeviceType>
 double PairLJCharmmCoulCharmmKokkos<DeviceType>::init_one(int i, int j)
 {
   double cutone = PairLJCharmmCoulCharmm::init_one(i,j);
   double cut_ljsqm = cut_ljsq;
   double cut_coulsqm = cut_coulsq;
 
   k_params.h_view(i,j).lj1 = lj1[i][j];
   k_params.h_view(i,j).lj2 = lj2[i][j];
   k_params.h_view(i,j).lj3 = lj3[i][j];
   k_params.h_view(i,j).lj4 = lj4[i][j];
   //k_params.h_view(i,j).offset = offset[i][j];
   k_params.h_view(i,j).cut_ljsq = cut_ljsqm;
   k_params.h_view(i,j).cut_coulsq = cut_coulsqm;
 
   k_params.h_view(j,i) = k_params.h_view(i,j);
   if(i<MAX_TYPES_STACKPARAMS+1 && j<MAX_TYPES_STACKPARAMS+1) {
     m_params[i][j] = m_params[j][i] = k_params.h_view(i,j);
     m_cutsq[j][i] = m_cutsq[i][j] = cutone*cutone;
     m_cut_ljsq[j][i] = m_cut_ljsq[i][j] = cut_ljsqm;
     m_cut_coulsq[j][i] = m_cut_coulsq[i][j] = cut_coulsqm;
   }
 
   k_cutsq.h_view(i,j) = cutone*cutone;
   k_cutsq.h_view(j,i) = k_cutsq.h_view(i,j);
   k_cutsq.template modify<LMPHostType>();
   k_cut_ljsq.h_view(i,j) = cut_ljsqm;
   k_cut_ljsq.h_view(j,i) = k_cut_ljsq.h_view(i,j);
   k_cut_ljsq.template modify<LMPHostType>();
   k_cut_coulsq.h_view(i,j) = cut_coulsqm;
   k_cut_coulsq.h_view(j,i) = k_cut_coulsq.h_view(i,j);
   k_cut_coulsq.template modify<LMPHostType>();
   k_params.template modify<LMPHostType>();
 
   return cutone;
 }
 
 
 
 namespace LAMMPS_NS {
 template class PairLJCharmmCoulCharmmKokkos<LMPDeviceType>;
 #ifdef KOKKOS_HAVE_CUDA
 template class PairLJCharmmCoulCharmmKokkos<LMPHostType>;
 #endif
 }
 
diff --git a/src/MOLECULE/pair_lj_charmm_coul_charmm.cpp b/src/MOLECULE/pair_lj_charmm_coul_charmm.cpp
index 0d08a672a..f12bc8f3b 100644
--- a/src/MOLECULE/pair_lj_charmm_coul_charmm.cpp
+++ b/src/MOLECULE/pair_lj_charmm_coul_charmm.cpp
@@ -1,532 +1,528 @@
 /* ----------------------------------------------------------------------
    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: Paul Crozier (SNL)
 ------------------------------------------------------------------------- */
 
 #include <math.h>
 #include <stdio.h>
 #include <stdlib.h>
 #include <string.h>
 #include "pair_lj_charmm_coul_charmm.h"
 #include "atom.h"
 #include "comm.h"
 #include "force.h"
 #include "neighbor.h"
 #include "neigh_list.h"
 #include "memory.h"
 #include "error.h"
 
 using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
 PairLJCharmmCoulCharmm::PairLJCharmmCoulCharmm(LAMMPS *lmp) : Pair(lmp)
 {
   implicit = 0;
   mix_flag = ARITHMETIC;
   writedata = 1;
 }
 
 /* ---------------------------------------------------------------------- */
 
 PairLJCharmmCoulCharmm::~PairLJCharmmCoulCharmm()
 {
   if (!copymode) {
    if (allocated) {
     memory->destroy(setflag);
     memory->destroy(cutsq);
 
     memory->destroy(epsilon);
     memory->destroy(sigma);
     memory->destroy(eps14);
     memory->destroy(sigma14);
     memory->destroy(lj1);
     memory->destroy(lj2);
     memory->destroy(lj3);
     memory->destroy(lj4);
     memory->destroy(lj14_1);
     memory->destroy(lj14_2);
     memory->destroy(lj14_3);
     memory->destroy(lj14_4);
    }
   }
 }
 
 /* ---------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::compute(int eflag, int vflag)
 {
   int i,j,ii,jj,inum,jnum,itype,jtype;
   double qtmp,xtmp,ytmp,ztmp,delx,dely,delz,evdwl,ecoul,fpair;
   double rsq,r2inv,r6inv,forcecoul,forcelj,factor_coul,factor_lj;
   double philj,switch1,switch2;
   int *ilist,*jlist,*numneigh,**firstneigh;
 
   evdwl = ecoul = 0.0;
   if (eflag || vflag) ev_setup(eflag,vflag);
   else evflag = vflag_fdotr = 0;
 
   double **x = atom->x;
   double **f = atom->f;
   double *q = atom->q;
   int *type = atom->type;
   int nlocal = atom->nlocal;
   double *special_coul = force->special_coul;
   double *special_lj = force->special_lj;
   int newton_pair = force->newton_pair;
   double qqrd2e = force->qqrd2e;
 
   inum = list->inum;
   ilist = list->ilist;
   numneigh = list->numneigh;
   firstneigh = list->firstneigh;
 
   // loop over neighbors of my atoms
 
   for (ii = 0; ii < inum; ii++) {
     i = ilist[ii];
     qtmp = q[i];
     xtmp = x[i][0];
     ytmp = x[i][1];
     ztmp = x[i][2];
     itype = type[i];
     jlist = firstneigh[i];
     jnum = numneigh[i];
 
     for (jj = 0; jj < jnum; jj++) {
       j = jlist[jj];
       factor_lj = special_lj[sbmask(j)];
       factor_coul = special_coul[sbmask(j)];
       j &= NEIGHMASK;
       jtype = type[j];
 
       delx = xtmp - x[j][0];
       dely = ytmp - x[j][1];
       delz = ztmp - x[j][2];
       rsq = delx*delx + dely*dely + delz*delz;
 
       if (rsq < cut_bothsq) {
         r2inv = 1.0/rsq;
 
         if (rsq < cut_coulsq) {
           forcecoul = qqrd2e * qtmp*q[j]*sqrt(r2inv);
           if (rsq > cut_coul_innersq) {
             switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
               (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) / denom_coul;
-            switch2 = 12.0*rsq * (cut_coulsq-rsq) *
-              (rsq-cut_coul_innersq) / denom_coul;
-            forcecoul *= switch1 + switch2;
+            forcecoul *= switch1;
           }
         } else forcecoul = 0.0;
 
         if (rsq < cut_ljsq) {
           r6inv = r2inv*r2inv*r2inv;
           forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
           if (rsq > cut_lj_innersq) {
             switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
               (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
             switch2 = 12.0*rsq * (cut_ljsq-rsq) *
               (rsq-cut_lj_innersq) / denom_lj;
             philj = r6inv * (lj3[itype][jtype]*r6inv - lj4[itype][jtype]);
             forcelj = forcelj*switch1 + philj*switch2;
           }
         } else forcelj = 0.0;
 
         fpair = (factor_coul*forcecoul + factor_lj*forcelj) * r2inv;
 
         f[i][0] += delx*fpair;
         f[i][1] += dely*fpair;
         f[i][2] += delz*fpair;
         if (newton_pair || j < nlocal) {
           f[j][0] -= delx*fpair;
           f[j][1] -= dely*fpair;
           f[j][2] -= delz*fpair;
         }
 
         if (eflag) {
           if (rsq < cut_coulsq) {
             ecoul = qqrd2e * qtmp*q[j]*sqrt(r2inv);
             if (rsq > cut_coul_innersq) {
               switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
                 (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) /
                 denom_coul;
               ecoul *= switch1;
             }
             ecoul *= factor_coul;
           } else ecoul = 0.0;
           if (rsq < cut_ljsq) {
             evdwl = r6inv*(lj3[itype][jtype]*r6inv-lj4[itype][jtype]);
             if (rsq > cut_lj_innersq) {
               switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
                 (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
               evdwl *= switch1;
             }
             evdwl *= factor_lj;
           } else evdwl = 0.0;
         }
 
         if (evflag) ev_tally(i,j,nlocal,newton_pair,
                              evdwl,ecoul,fpair,delx,dely,delz);
       }
     }
   }
 
   if (vflag_fdotr) virial_fdotr_compute();
 }
 
 /* ----------------------------------------------------------------------
    allocate all arrays
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::allocate()
 {
   allocated = 1;
   int n = atom->ntypes;
 
   memory->create(setflag,n+1,n+1,"pair:setflag");
   for (int i = 1; i <= n; i++)
     for (int j = i; j <= n; j++)
       setflag[i][j] = 0;
 
   memory->create(cutsq,n+1,n+1,"pair:cutsq");
 
   memory->create(epsilon,n+1,n+1,"pair:epsilon");
   memory->create(sigma,n+1,n+1,"pair:sigma");
   memory->create(eps14,n+1,n+1,"pair:eps14");
   memory->create(sigma14,n+1,n+1,"pair:sigma14");
   memory->create(lj1,n+1,n+1,"pair:lj1");
   memory->create(lj2,n+1,n+1,"pair:lj2");
   memory->create(lj3,n+1,n+1,"pair:lj3");
   memory->create(lj4,n+1,n+1,"pair:lj4");
   memory->create(lj14_1,n+1,n+1,"pair:lj14_1");
   memory->create(lj14_2,n+1,n+1,"pair:lj14_2");
   memory->create(lj14_3,n+1,n+1,"pair:lj14_3");
   memory->create(lj14_4,n+1,n+1,"pair:lj14_4");
 }
 
 /* ----------------------------------------------------------------------
    global settings
    unlike other pair styles,
      there are no individual pair settings that these override
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::settings(int narg, char **arg)
 {
   if (narg != 2 && narg != 4)
     error->all(FLERR,"Illegal pair_style command");
 
   cut_lj_inner = force->numeric(FLERR,arg[0]);
   cut_lj = force->numeric(FLERR,arg[1]);
   if (narg == 2) {
     cut_coul_inner = cut_lj_inner;
     cut_coul = cut_lj;
   } else {
     cut_coul_inner = force->numeric(FLERR,arg[2]);
     cut_coul = force->numeric(FLERR,arg[3]);
   }
 }
 
 /* ----------------------------------------------------------------------
    set coeffs for one or more type pairs
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::coeff(int narg, char **arg)
 {
   if (narg != 4 && narg != 6)
     error->all(FLERR,"Incorrect args for pair coefficients");
   if (!allocated) allocate();
 
   int ilo,ihi,jlo,jhi;
   force->bounds(arg[0],atom->ntypes,ilo,ihi);
   force->bounds(arg[1],atom->ntypes,jlo,jhi);
 
   double epsilon_one = force->numeric(FLERR,arg[2]);
   double sigma_one = force->numeric(FLERR,arg[3]);
   double eps14_one = epsilon_one;
   double sigma14_one = sigma_one;
   if (narg == 6) {
     eps14_one = force->numeric(FLERR,arg[4]);
     sigma14_one = force->numeric(FLERR,arg[5]);
   }
 
   int count = 0;
   for (int i = ilo; i <= ihi; i++) {
     for (int j = MAX(jlo,i); j <= jhi; j++) {
       epsilon[i][j] = epsilon_one;
       sigma[i][j] = sigma_one;
       eps14[i][j] = eps14_one;
       sigma14[i][j] = sigma14_one;
       setflag[i][j] = 1;
       count++;
     }
   }
 
   if (count == 0) error->all(FLERR,"Incorrect args for pair coefficients");
 }
 
 /* ----------------------------------------------------------------------
    init specific to this pair style
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::init_style()
 {
   if (!atom->q_flag)
     error->all(FLERR,
                "Pair style lj/charmm/coul/charmm requires atom attribute q");
 
   neighbor->request(this,instance_me);
 
   // require cut_lj_inner < cut_lj, cut_coul_inner < cut_coul
 
   if (cut_lj_inner >= cut_lj || cut_coul_inner >= cut_coul)
     error->all(FLERR,"Pair inner cutoff >= Pair outer cutoff");
 
   cut_lj_innersq = cut_lj_inner * cut_lj_inner;
   cut_ljsq = cut_lj * cut_lj;
   cut_coul_innersq = cut_coul_inner * cut_coul_inner;
   cut_coulsq = cut_coul * cut_coul;
   cut_bothsq = MAX(cut_ljsq,cut_coulsq);
 
   denom_lj = (cut_ljsq-cut_lj_innersq) * (cut_ljsq-cut_lj_innersq) *
     (cut_ljsq-cut_lj_innersq);
   denom_coul = (cut_coulsq-cut_coul_innersq) * (cut_coulsq-cut_coul_innersq) *
     (cut_coulsq-cut_coul_innersq);
 }
 
 /* ----------------------------------------------------------------------
    init for one type pair i,j and corresponding j,i
 ------------------------------------------------------------------------- */
 
 double PairLJCharmmCoulCharmm::init_one(int i, int j)
 {
   if (setflag[i][j] == 0) {
     epsilon[i][j] = mix_energy(epsilon[i][i],epsilon[j][j],
                                sigma[i][i],sigma[j][j]);
     sigma[i][j] = mix_distance(sigma[i][i],sigma[j][j]);
     eps14[i][j] = mix_energy(eps14[i][i],eps14[j][j],
                                sigma14[i][i],sigma14[j][j]);
     sigma14[i][j] = mix_distance(sigma14[i][i],sigma14[j][j]);
   }
 
   double cut = MAX(cut_lj,cut_coul);
 
   lj1[i][j] = 48.0 * epsilon[i][j] * pow(sigma[i][j],12.0);
   lj2[i][j] = 24.0 * epsilon[i][j] * pow(sigma[i][j],6.0);
   lj3[i][j] = 4.0 * epsilon[i][j] * pow(sigma[i][j],12.0);
   lj4[i][j] = 4.0 * epsilon[i][j] * pow(sigma[i][j],6.0);
   lj14_1[i][j] = 48.0 * eps14[i][j] * pow(sigma14[i][j],12.0);
   lj14_2[i][j] = 24.0 * eps14[i][j] * pow(sigma14[i][j],6.0);
   lj14_3[i][j] = 4.0 * eps14[i][j] * pow(sigma14[i][j],12.0);
   lj14_4[i][j] = 4.0 * eps14[i][j] * pow(sigma14[i][j],6.0);
 
   lj1[j][i] = lj1[i][j];
   lj2[j][i] = lj2[i][j];
   lj3[j][i] = lj3[i][j];
   lj4[j][i] = lj4[i][j];
   lj14_1[j][i] = lj14_1[i][j];
   lj14_2[j][i] = lj14_2[i][j];
   lj14_3[j][i] = lj14_3[i][j];
   lj14_4[j][i] = lj14_4[i][j];
 
   return cut;
 }
 
 /* ----------------------------------------------------------------------
   proc 0 writes to restart file
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::write_restart(FILE *fp)
 {
   write_restart_settings(fp);
 
   int i,j;
   for (i = 1; i <= atom->ntypes; i++)
     for (j = i; j <= atom->ntypes; j++) {
       fwrite(&setflag[i][j],sizeof(int),1,fp);
       if (setflag[i][j]) {
         fwrite(&epsilon[i][j],sizeof(double),1,fp);
         fwrite(&sigma[i][j],sizeof(double),1,fp);
         fwrite(&eps14[i][j],sizeof(double),1,fp);
         fwrite(&sigma14[i][j],sizeof(double),1,fp);
       }
     }
 }
 
 /* ----------------------------------------------------------------------
   proc 0 reads from restart file, bcasts
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::read_restart(FILE *fp)
 {
   read_restart_settings(fp);
 
   allocate();
 
   int i,j;
   int me = comm->me;
   for (i = 1; i <= atom->ntypes; i++)
     for (j = i; j <= atom->ntypes; j++) {
       if (me == 0) fread(&setflag[i][j],sizeof(int),1,fp);
       MPI_Bcast(&setflag[i][j],1,MPI_INT,0,world);
       if (setflag[i][j]) {
         if (me == 0) {
           fread(&epsilon[i][j],sizeof(double),1,fp);
           fread(&sigma[i][j],sizeof(double),1,fp);
           fread(&eps14[i][j],sizeof(double),1,fp);
           fread(&sigma14[i][j],sizeof(double),1,fp);
         }
         MPI_Bcast(&epsilon[i][j],1,MPI_DOUBLE,0,world);
         MPI_Bcast(&sigma[i][j],1,MPI_DOUBLE,0,world);
         MPI_Bcast(&eps14[i][j],1,MPI_DOUBLE,0,world);
         MPI_Bcast(&sigma14[i][j],1,MPI_DOUBLE,0,world);
       }
     }
 }
 
 /* ----------------------------------------------------------------------
    proc 0 writes to data file
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::write_data(FILE *fp)
 {
   for (int i = 1; i <= atom->ntypes; i++)
     fprintf(fp,"%d %g %g %g %g\n",
             i,epsilon[i][i],sigma[i][i],eps14[i][i],sigma14[i][i]);
 }
 
 /* ----------------------------------------------------------------------
    proc 0 writes all pairs to data file
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::write_data_all(FILE *fp)
 {
   for (int i = 1; i <= atom->ntypes; i++)
     for (int j = i; j <= atom->ntypes; j++)
       fprintf(fp,"%d %d %g %g %g %g\n",i,j,
               epsilon[i][j],sigma[i][j],eps14[i][j],sigma14[i][j]);
 }
 
 /* ----------------------------------------------------------------------
   proc 0 writes to restart file
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::write_restart_settings(FILE *fp)
 {
   fwrite(&cut_lj_inner,sizeof(double),1,fp);
   fwrite(&cut_lj,sizeof(double),1,fp);
   fwrite(&cut_coul_inner,sizeof(double),1,fp);
   fwrite(&cut_coul,sizeof(double),1,fp);
   fwrite(&offset_flag,sizeof(int),1,fp);
   fwrite(&mix_flag,sizeof(int),1,fp);
 }
 
 /* ----------------------------------------------------------------------
   proc 0 reads from restart file, bcasts
 ------------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmm::read_restart_settings(FILE *fp)
 {
   if (comm->me == 0) {
     fread(&cut_lj_inner,sizeof(double),1,fp);
     fread(&cut_lj,sizeof(double),1,fp);
     fread(&cut_coul_inner,sizeof(double),1,fp);
     fread(&cut_coul,sizeof(double),1,fp);
     fread(&offset_flag,sizeof(int),1,fp);
     fread(&mix_flag,sizeof(int),1,fp);
   }
   MPI_Bcast(&cut_lj_inner,1,MPI_DOUBLE,0,world);
   MPI_Bcast(&cut_lj,1,MPI_DOUBLE,0,world);
   MPI_Bcast(&cut_coul_inner,1,MPI_DOUBLE,0,world);
   MPI_Bcast(&cut_coul,1,MPI_DOUBLE,0,world);
   MPI_Bcast(&offset_flag,1,MPI_INT,0,world);
   MPI_Bcast(&mix_flag,1,MPI_INT,0,world);
 }
 
 /* ---------------------------------------------------------------------- */
 
 double PairLJCharmmCoulCharmm::single(int i, int j, int itype, int jtype,
                                       double rsq,
                                       double factor_coul, double factor_lj,
                                       double &fforce)
 {
   double r2inv,r6inv,forcecoul,forcelj,phicoul,philj;
   double switch1,switch2;
 
   r2inv = 1.0/rsq;
   if (rsq < cut_coulsq) {
     forcecoul = force->qqrd2e * atom->q[i]*atom->q[j]*sqrt(r2inv);
     if (rsq > cut_coul_innersq) {
       switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
         (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) / denom_coul;
-      switch2 = 12.0*rsq * (cut_coulsq-rsq) *
-        (rsq-cut_coul_innersq) / denom_coul;
-      forcecoul *= switch1 + switch2;
+      forcecoul *= switch1;
     }
   } else forcecoul = 0.0;
 
   if (rsq < cut_ljsq) {
     r6inv = r2inv*r2inv*r2inv;
     forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
     if (rsq > cut_lj_innersq) {
       switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
         (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
       switch2 = 12.0*rsq * (cut_ljsq-rsq) *
         (rsq-cut_lj_innersq) / denom_lj;
       philj = r6inv * (lj3[itype][jtype]*r6inv - lj4[itype][jtype]);
       forcelj = forcelj*switch1 + philj*switch2;
     }
   } else forcelj = 0.0;
 
   fforce = (factor_coul*forcecoul + factor_lj*forcelj) * r2inv;
 
   double eng = 0.0;
   if (rsq < cut_coulsq) {
     phicoul = force->qqrd2e * atom->q[i]*atom->q[j]*sqrt(r2inv);
     if (rsq > cut_coul_innersq) {
       switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
         (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) /
         denom_coul;
       phicoul *= switch1;
     }
     eng += factor_coul*phicoul;
   }
   if (rsq < cut_ljsq) {
     philj = r6inv*(lj3[itype][jtype]*r6inv-lj4[itype][jtype]);
     if (rsq > cut_lj_innersq) {
       switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
         (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) / denom_lj;
       philj *= switch1;
     }
     eng += factor_lj*philj;
   }
 
   return eng;
 }
 
 /* ---------------------------------------------------------------------- */
 
 void *PairLJCharmmCoulCharmm::extract(const char *str, int &dim)
 {
   dim = 2;
   if (strcmp(str,"lj14_1") == 0) return (void *) lj14_1;
   if (strcmp(str,"lj14_2") == 0) return (void *) lj14_2;
   if (strcmp(str,"lj14_3") == 0) return (void *) lj14_3;
   if (strcmp(str,"lj14_4") == 0) return (void *) lj14_4;
 
   dim = 0;
   if (strcmp(str,"implicit") == 0) return (void *) &implicit;
 
   return NULL;
 }
diff --git a/src/USER-OMP/pair_lj_charmm_coul_charmm_omp.cpp b/src/USER-OMP/pair_lj_charmm_coul_charmm_omp.cpp
index 880fa4006..9fd41e61d 100644
--- a/src/USER-OMP/pair_lj_charmm_coul_charmm_omp.cpp
+++ b/src/USER-OMP/pair_lj_charmm_coul_charmm_omp.cpp
@@ -1,212 +1,210 @@
 /* ----------------------------------------------------------------------
    LAMMPS - Large-scale Atomic/Molecular Massively Parallel Simulator
    http://lammps.sandia.gov, Sandia National Laboratories
    Steve Plimpton, sjplimp@sandia.gov
 
    This software is distributed under the GNU General Public License.
 
    See the README file in the top-level LAMMPS directory.
 ------------------------------------------------------------------------- */
 
 /* ----------------------------------------------------------------------
    Contributing author: Axel Kohlmeyer (Temple U)
 ------------------------------------------------------------------------- */
 
 #include <math.h>
 #include "pair_lj_charmm_coul_charmm_omp.h"
 #include "atom.h"
 #include "comm.h"
 #include "force.h"
 #include "neighbor.h"
 #include "neigh_list.h"
 
 #include "suffix.h"
 using namespace LAMMPS_NS;
 
 /* ---------------------------------------------------------------------- */
 
 PairLJCharmmCoulCharmmOMP::PairLJCharmmCoulCharmmOMP(LAMMPS *lmp) :
   PairLJCharmmCoulCharmm(lmp), ThrOMP(lmp, THR_PAIR)
 {
   suffix_flag |= Suffix::OMP;
   respa_enable = 0;
 }
 
 /* ---------------------------------------------------------------------- */
 
 void PairLJCharmmCoulCharmmOMP::compute(int eflag, int vflag)
 {
   if (eflag || vflag) {
     ev_setup(eflag,vflag);
   } else evflag = vflag_fdotr = 0;
 
   const int nall = atom->nlocal + atom->nghost;
   const int nthreads = comm->nthreads;
   const int inum = list->inum;
 
 #if defined(_OPENMP)
 #pragma omp parallel default(none) shared(eflag,vflag)
 #endif
   {
     int ifrom, ito, tid;
 
     loop_setup_thr(ifrom, ito, tid, inum, nthreads);
     ThrData *thr = fix->get_thr(tid);
     thr->timer(Timer::START);
     ev_setup_thr(eflag, vflag, nall, eatom, vatom, thr);
 
     if (evflag) {
       if (eflag) {
         if (force->newton_pair) eval<1,1,1>(ifrom, ito, thr);
         else eval<1,1,0>(ifrom, ito, thr);
       } else {
         if (force->newton_pair) eval<1,0,1>(ifrom, ito, thr);
         else eval<1,0,0>(ifrom, ito, thr);
       }
     } else {
       if (force->newton_pair) eval<0,0,1>(ifrom, ito, thr);
       else eval<0,0,0>(ifrom, ito, thr);
     }
 
     thr->timer(Timer::PAIR);
     reduce_thr(this, eflag, vflag, thr);
   } // end of omp parallel region
 }
 
 /* ---------------------------------------------------------------------- */
 
 template <int EVFLAG, int EFLAG, int NEWTON_PAIR>
 void PairLJCharmmCoulCharmmOMP::eval(int iifrom, int iito, ThrData * const thr)
 {
   int i,j,ii,jj,jnum,itype,jtype;
   double qtmp,xtmp,ytmp,ztmp,delx,dely,delz,evdwl,ecoul,fpair;
   double rsq,r2inv,r6inv,forcecoul,forcelj,factor_coul,factor_lj;
   double philj,switch1,switch2;
   int *ilist,*jlist,*numneigh,**firstneigh;
 
   evdwl = ecoul = 0.0;
 
   const dbl3_t * _noalias const x = (dbl3_t *) atom->x[0];
   dbl3_t * _noalias const f = (dbl3_t *) thr->get_f()[0];
   const double * _noalias const q = atom->q;
   const int * _noalias const type = atom->type;
   const int nlocal = atom->nlocal;
   const double * _noalias const special_coul = force->special_coul;
   const double * _noalias const special_lj = force->special_lj;
   const double qqrd2e = force->qqrd2e;
   double fxtmp,fytmp,fztmp;
 
   ilist = list->ilist;
   numneigh = list->numneigh;
   firstneigh = list->firstneigh;
   const double invdenom_coul = (denom_coul != 0.0) ? 1.0/denom_coul : 0.0;
   const double invdenom_lj   = (denom_lj   != 0.0) ? 1.0/denom_lj   : 0.0;
 
   // loop over neighbors of my atoms
 
   for (ii = iifrom; ii < iito; ++ii) {
 
     i = ilist[ii];
     qtmp = q[i];
     xtmp = x[i].x;
     ytmp = x[i].y;
     ztmp = x[i].z;
     itype = type[i];
     jlist = firstneigh[i];
     jnum = numneigh[i];
     fxtmp=fytmp=fztmp=0.0;
 
     for (jj = 0; jj < jnum; jj++) {
       j = jlist[jj];
       factor_lj = special_lj[sbmask(j)];
       factor_coul = special_coul[sbmask(j)];
       j &= NEIGHMASK;
 
       delx = xtmp - x[j].x;
       dely = ytmp - x[j].y;
       delz = ztmp - x[j].z;
       rsq = delx*delx + dely*dely + delz*delz;
       jtype = type[j];
 
       if (rsq < cut_bothsq) {
         r2inv = 1.0/rsq;
 
         if (rsq < cut_coulsq) {
           forcecoul = qqrd2e * qtmp*q[j]*sqrt(r2inv);
           if (rsq > cut_coul_innersq) {
             switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
               (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) * invdenom_coul;
-            switch2 = 12.0*rsq * (cut_coulsq-rsq) *
-              (rsq-cut_coul_innersq) * invdenom_coul;
-            forcecoul *= switch1 + switch2;
+            forcecoul *= switch1;
           }
           forcecoul *= factor_coul;
         } else forcecoul = 0.0;
 
         if (rsq < cut_ljsq) {
           r6inv = r2inv*r2inv*r2inv;
           jtype = type[j];
           forcelj = r6inv * (lj1[itype][jtype]*r6inv - lj2[itype][jtype]);
           if (rsq > cut_lj_innersq) {
             switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
               (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) * invdenom_lj;
             switch2 = 12.0*rsq * (cut_ljsq-rsq) *
               (rsq-cut_lj_innersq) * invdenom_lj;
             philj = r6inv * (lj3[itype][jtype]*r6inv - lj4[itype][jtype]);
             forcelj = forcelj*switch1 + philj*switch2;
           }
           forcelj *= factor_lj;
         } else forcelj = 0.0;
 
         fpair = (forcecoul + forcelj) * 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 (EFLAG) {
           if (rsq < cut_coulsq) {
             ecoul = qqrd2e * qtmp*q[j]*sqrt(r2inv);
             if (rsq > cut_coul_innersq) {
               switch1 = (cut_coulsq-rsq) * (cut_coulsq-rsq) *
                 (cut_coulsq + 2.0*rsq - 3.0*cut_coul_innersq) *
                 invdenom_coul;
               ecoul *= switch1;
             }
             ecoul *= factor_coul;
           } else ecoul = 0.0;
           if (rsq < cut_ljsq) {
             evdwl = r6inv*(lj3[itype][jtype]*r6inv-lj4[itype][jtype]);
             if (rsq > cut_lj_innersq) {
               switch1 = (cut_ljsq-rsq) * (cut_ljsq-rsq) *
                 (cut_ljsq + 2.0*rsq - 3.0*cut_lj_innersq) * invdenom_lj;
               evdwl *= switch1;
             }
             evdwl *= factor_lj;
           } else evdwl = 0.0;
         }
 
         if (EVFLAG) ev_tally_thr(this, i,j,nlocal,NEWTON_PAIR,
                                  evdwl,ecoul,fpair,delx,dely,delz,thr);
       }
     }
     f[i].x += fxtmp;
     f[i].y += fytmp;
     f[i].z += fztmp;
   }
 }
 
 /* ---------------------------------------------------------------------- */
 
 double PairLJCharmmCoulCharmmOMP::memory_usage()
 {
   double bytes = memory_usage_thr();
   bytes += PairLJCharmmCoulCharmm::memory_usage();
 
   return bytes;
 }