diff --git a/lib/gpu/lal_lj_coul_soft.cpp b/lib/gpu/lal_lj_coul_soft.cpp index 96bcf41aa3..9ee6486817 100644 --- a/lib/gpu/lal_lj_coul_soft.cpp +++ b/lib/gpu/lal_lj_coul_soft.cpp @@ -56,7 +56,7 @@ int LJCoulSoftT::init(const int ntypes, double *host_special_coul, const double qqrd2e) { int success; success=this->init_atomic(nlocal,nall,max_nbors,maxspecial,cell_size,gpu_split, - _screen,lj_coul_soft,"lj_coul_soft"); + _screen,lj_coul_soft,"k_lj_coul_soft"); if (success!=0) return success; diff --git a/lib/gpu/lal_lj_coul_soft.cu b/lib/gpu/lal_lj_coul_soft.cu index 74da36e2c1..1fc564bde6 100644 --- a/lib/gpu/lal_lj_coul_soft.cu +++ b/lib/gpu/lal_lj_coul_soft.cu @@ -225,7 +225,7 @@ __kernel void k_lj_coul_soft_fast(const __global numtyp4 *restrict x_, numtyp forcecoul, force_lj, force; numtyp r4sig6, denlj, denc; - if (rsq < lj1[mtype].z) { + if (rsq < lj1[mtype].z) { // cut_ljsq[itype][jtype] r4sig6 = rsq*rsq / lj1[mtype].y; denlj = lj3[mtype].x + rsq*r4sig6; force_lj = lj1[mtype].x * lj3[mtype].w * @@ -234,7 +234,7 @@ __kernel void k_lj_coul_soft_fast(const __global numtyp4 *restrict x_, } else force_lj = (numtyp)0.0; - if (rsq < lj1[mtype].w) { + if (rsq < lj1[mtype].w) { // cut_coulsq[itype][jtype] fetch(forcecoul,j,q_tex); denc = sqrt(lj3[mtype].y + rsq); forcecoul *= qqrd2e * lj1[mtype].x * qtmp / (denc*denc*denc); diff --git a/lib/gpu/lal_lj_coul_soft_ext.cpp b/lib/gpu/lal_lj_coul_soft_ext.cpp index 8c10db14f6..02d367b3c7 100644 --- a/lib/gpu/lal_lj_coul_soft_ext.cpp +++ b/lib/gpu/lal_lj_coul_soft_ext.cpp @@ -43,7 +43,7 @@ int ljcs_gpu_init(const int ntypes, double **cutsq, double **host_lj1, int gpu_rank=LJCSMF.device->gpu_rank(); int procs_per_gpu=LJCSMF.device->procs_per_gpu(); - LJCSMF.device->init_message(screen,"lj/cut/coul/cut",first_gpu,last_gpu); + LJCSMF.device->init_message(screen,"lj/cut/coul/cut/soft",first_gpu,last_gpu); bool message=false; if (LJCSMF.device->replica_me()==0 && screen) diff --git a/src/GPU/pair_lj_cut_coul_cut_soft_gpu.cpp b/src/GPU/pair_lj_cut_coul_cut_soft_gpu.cpp index d65e3bd4cf..cfde3ab632 100644 --- a/src/GPU/pair_lj_cut_coul_cut_soft_gpu.cpp +++ b/src/GPU/pair_lj_cut_coul_cut_soft_gpu.cpp @@ -110,6 +110,8 @@ void PairLJCutCoulCutSoftGPU::compute(int eflag, int vflag) } if (!success) error->one(FLERR, "Insufficient memory on accelerator"); + if (atom->molecular != Atom::ATOMIC && neighbor->ago == 0) + neighbor->build_topology(); if (host_start < inum) { cpu_time = platform::walltime(); cpu_compute(host_start, inum, eflag, vflag, ilist, numneigh, firstneigh); @@ -227,13 +229,11 @@ void PairLJCutCoulCutSoftGPU::cpu_compute(int start, int inum, int eflag, int /* f[i][1] += dely * fpair; f[i][2] += delz * fpair; - if (eflag) { if (rsq < cut_coulsq[itype][jtype]) ecoul = factor_coul * qqrd2e * lj1[itype][jtype] * qtmp*q[j] / denc; else ecoul = 0.0; - if (rsq < cut_ljsq[itype][jtype]) { evdwl = lj1[itype][jtype] * 4.0 * epsilon[itype][jtype] * (1.0/(denlj*denlj) - 1.0/denlj) - offset[itype][jtype];