git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@7886 f3b2605a-c512-4ea7-a41b-209d697bcdaa

src/GPU/gpu_extra.h

@@ -29,7 +29,8 @@ namespace GPU_EXTRA {
     MPI_Allreduce(&error_flag, &all_success, 1, MPI_INT, MPI_MIN, world);
     if (all_success != 0) {
       if (all_success == -1)
-	error->all(FLERR,"Accelerated style in input script but no fix gpu"); 
+	error->all(FLERR,
+		   "The package gpu command is required for gpu styles"); 
       else if (all_success == -2)
 	error->all(FLERR,
 		   "Could not find/initialize a specified accelerator device");
@@ -64,44 +65,43 @@ namespace GPU_EXTRA {
 
 /* ERROR/WARNING messages:
 
-E: Accelerated style in input script but no fix gpu
+E: The package gpu command is required for gpu styles
 
-UNDOCUMENTED
+Self-explanatory.
 
 E: Could not find/initialize a specified accelerator device
 
-UNDOCUMENTED
+Could not initialize at least one of the devices specified for the gpu package
 
 E: Insufficient memory on accelerator
 
-UNDOCUMENTED
+There is insufficient memory on one of the devices specified for the gpu 
+package
 
 E: GPU library not compiled for this accelerator
 
-UNDOCUMENTED
+Self-explanatory.
 
 E: Double precision is not supported on this accelerator
 
-UNDOCUMENTED
+Self-explanatory
 
 E: Unable to initialize accelerator for use
 
-UNDOCUMENTED
+There was a problem initializing an accelerator for the gpu package
 
 E: Accelerator sharing is not currently supported on system
 
-UNDOCUMENTED
+Multiple MPI processes cannot share the accelerator on your system. For NVIDIA
+GPUs, see the nvidia-smi command to change this setting.
 
 E: GPU particle split must be set to 1 for this pair style.
 
-UNDOCUMENTED
+For this pair style, you cannot run part of the force calculation on the host.
+See the package command.
 
 E: Unknown error in GPU library
 
-UNDOCUMENTED
+Self-explanatory.
-
-E: The package gpu command is required for gpu styles
-
-UNDOCUMENTED
 
 */

src/GPU/pair_buck_coul_cut_gpu.cpp

@@ -115,7 +115,7 @@ void PairBuckCoulCutGPU::compute(int eflag, int vflag)
 		      atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_buck_coul_cut_gpu.h

@@ -47,10 +47,10 @@ class PairBuckCoulCutGPU : public PairBuckCoulCut {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with buck/coul/cut/gpu pair style
 

src/GPU/pair_buck_coul_long_gpu.cpp

@@ -124,7 +124,7 @@ void PairBuckCoulLongGPU::compute(int eflag, int vflag)
 		       atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_buck_coul_long_gpu.h

@@ -47,10 +47,10 @@ class PairBuckCoulLongGPU : public PairBuckCoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style buck/coul/long/gpu requires atom attribute q
 

src/GPU/pair_buck_gpu.cpp

@@ -108,7 +108,7 @@ void PairBuckGPU::compute(int eflag, int vflag)
 		     vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_buck_gpu.h

@@ -47,10 +47,10 @@ class PairBuckGPU : public PairBuck {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with buck/gpu pair style
 

src/GPU/pair_coul_long_gpu.cpp

@@ -120,7 +120,7 @@ void PairCoulLongGPU::compute(int eflag, int vflag)
 		   atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_coul_long_gpu.h

@@ -47,10 +47,10 @@ class PairCoulLongGPU : public PairCoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style coul/long/gpu requires atom attribute q
 

src/GPU/pair_eam_gpu.cpp

@@ -130,7 +130,7 @@ void PairEAMGPU::compute(int eflag, int vflag)
   }
     
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   // communicate derivative of embedding function
 
@@ -152,8 +152,6 @@ void PairEAMGPU::init_style()
   if (force->newton_pair) 
     error->all(FLERR,"Cannot use newton pair with eam/gpu pair style");
   
-  if (!allocated) error->all(FLERR,"Not allocate memory eam/gpu pair style");
-  
   // convert read-in file(s) to arrays and spline them
 
   file2array();

src/GPU/pair_eam_gpu.h

@@ -54,17 +54,13 @@ class PairEAMGPU : public PairEAM {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with eam/gpu pair style
 
 Self-explanatory.
 
-E: Not allocate memory eam/gpu pair style
-
-UNDOCUMENTED
-
 */

src/GPU/pair_gayberne_gpu.cpp

@@ -137,7 +137,7 @@ void PairGayBerneGPU::compute(int eflag, int vflag)
 			   cpu_time, success, quat);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start < inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_gayberne_gpu.h

@@ -53,10 +53,10 @@ E: Pair gayberne/gpu requires atom style ellipsoid
 
 Self-explanatory.
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with gayberne/gpu pair style
 

src/GPU/pair_lj96_cut_gpu.cpp

@@ -107,7 +107,7 @@ void PairLJ96CutGPU::compute(int eflag, int vflag)
 		     vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj96_cut_gpu.h

@@ -47,10 +47,10 @@ class PairLJ96CutGPU : public PairLJ96Cut {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with lj96/cut/gpu pair style
 

src/GPU/pair_lj_charmm_coul_long_gpu.cpp

@@ -126,7 +126,7 @@ void PairLJCharmmCoulLongGPU::compute(int eflag, int vflag)
 		     atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_charmm_coul_long_gpu.h

@@ -47,10 +47,10 @@ class PairLJCharmmCoulLongGPU : public PairLJCharmmCoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style lj/charmm/coul/long/gpu requires atom attribute q
 

src/GPU/pair_lj_class2_coul_long_gpu.cpp

@@ -123,7 +123,7 @@ void PairLJClass2CoulLongGPU::compute(int eflag, int vflag)
 		     atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_class2_coul_long_gpu.h

@@ -47,10 +47,10 @@ class PairLJClass2CoulLongGPU : public PairLJClass2CoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style lj/class2/coul/long/gpu requires atom attribute q
 

src/GPU/pair_lj_class2_gpu.cpp

@@ -106,7 +106,7 @@ void PairLJClass2GPU::compute(int eflag, int vflag)
 		     vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_class2_gpu.h

@@ -47,10 +47,10 @@ class PairLJClass2GPU : public PairLJClass2 {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with lj/class2/gpu pair style
 

src/GPU/pair_lj_cut_coul_cut_gpu.cpp

@@ -114,7 +114,7 @@ void PairLJCutCoulCutGPU::compute(int eflag, int vflag)
 		    atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_cut_coul_cut_gpu.h

@@ -47,10 +47,10 @@ class PairLJCutCoulCutGPU : public PairLJCutCoulCut {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style lj/cut/coul/cut/gpu requires atom attribute q
 

src/GPU/pair_lj_cut_coul_long_gpu.cpp

@@ -124,7 +124,7 @@ void PairLJCutCoulLongGPU::compute(int eflag, int vflag)
 		     atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_cut_coul_long_gpu.h

@@ -47,10 +47,10 @@ class PairLJCutCoulLongGPU : public PairLJCutCoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style lj/cut/coul/long/gpu requires atom attribute q
 

src/GPU/pair_lj_cut_gpu.cpp

@@ -107,7 +107,7 @@ void PairLJCutGPU::compute(int eflag, int vflag)
 		    vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_cut_gpu.h

@@ -47,10 +47,10 @@ class PairLJCutGPU : public PairLJCut {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with lj/cut/gpu pair style
 

src/GPU/pair_lj_expand_gpu.cpp

@@ -108,7 +108,7 @@ void PairLJExpandGPU::compute(int eflag, int vflag)
 		    vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_expand_gpu.h

@@ -47,10 +47,10 @@ class PairLJExpandGPU : public PairLJExpand {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with lj/expand/gpu pair style
 

src/GPU/pair_lj_sdk_coul_long_gpu.cpp

@@ -127,7 +127,7 @@ void PairLJSDKCoulLongGPU::compute(int eflag, int vflag)
 		     atom->nlocal, domain->boxlo, domain->prd);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_sdk_coul_long_gpu.h

@@ -49,10 +49,10 @@ class PairLJSDKCoulLongGPU : public PairLJSDKCoulLong {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Pair style lj/sdk/coul/long/gpu requires atom attribute q
 

src/GPU/pair_lj_sdk_gpu.cpp

@@ -111,7 +111,7 @@ void PairLJSDKGPU::compute(int eflag, int vflag)
 		    vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_lj_sdk_gpu.h

@@ -49,10 +49,10 @@ class PairLJSDKGPU : public PairLJSDK {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with lj/sdk/gpu pair style
 

src/GPU/pair_morse_gpu.cpp

@@ -106,7 +106,7 @@ void PairMorseGPU::compute(int eflag, int vflag)
 		    vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_morse_gpu.h

@@ -47,10 +47,10 @@ class PairMorseGPU : public PairMorse {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with morse/gpu pair style
 

src/GPU/pair_resquared_gpu.cpp

@@ -136,7 +136,7 @@ void PairRESquaredGPU::compute(int eflag, int vflag)
 			   cpu_time, success, quat);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start < inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_resquared_gpu.h

@@ -53,10 +53,10 @@ E: Pair resquared/gpu requires atom style ellipsoid
 
 Self-explanatory.
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with resquared/gpu pair style
 

src/GPU/pair_table_gpu.cpp

@@ -114,7 +114,7 @@ void PairTableGPU::compute(int eflag, int vflag)
 		      vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_table_gpu.h

@@ -47,10 +47,10 @@ class PairTableGPU : public PairTable {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with table/gpu pair style
 

src/GPU/pair_yukawa_gpu.cpp

@@ -108,7 +108,7 @@ void PairYukawaGPU::compute(int eflag, int vflag)
 		       vflag_atom, host_start, cpu_time, success);
   }
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
 
   if (host_start<inum) {
     cpu_time = MPI_Wtime();

src/GPU/pair_yukawa_gpu.h

@@ -47,10 +47,10 @@ class PairYukawaGPU : public PairYukawa {
 
 /* ERROR/WARNING messages:
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Cannot use newton pair with yukawa/gpu pair style
 

src/GPU/pppm_gpu.cpp

@@ -172,11 +172,21 @@ void PPPMGPU::compute(int eflag, int vflag)
 			     atom->q, domain->boxlo, delxinv, delyinv,
 			     delzinv);
   if (!success)
-    error->one(FLERR,"Out of memory on GPGPU");
+    error->one(FLERR,"Insufficient memory on accelerator");
   if (flag != 0)
     error->one(FLERR,"Out of range atoms - cannot compute PPPM");
 
-  int i;
+  // If need per-atom energies/virials, also do particle map on host
+  // concurrently with GPU calculations
+
+  if (evflag_atom) {
+    memory->destroy(part2grid);
+    nmax = atom->nmax;
+    memory->create(part2grid,nmax,3,"pppm:part2grid");
+    particle_map();
+  }
+
+  int i,j;
 
   // convert atoms from box to lamda coords
   
@@ -212,6 +222,29 @@ void PPPMGPU::compute(int eflag, int vflag)
   FFT_SCALAR qscale = force->qqrd2e * scale;
   PPPM_GPU_API(interp)(qscale);
 
+  // Compute per-atom energy/virial on host if requested
+
+  if (evflag_atom) {
+    fillbrick_peratom();
+    fieldforce_peratom();
+    double *q = atom->q;
+    int nlocal = atom->nlocal;
+
+    if (eflag_atom) {
+      for (i = 0; i < nlocal; i++) {
+	eatom[i] *= 0.5;
+        eatom[i] -= g_ewald*q[i]*q[i]/MY_PIS + MY_PI2*q[i]*qsum / 
+	  (g_ewald*g_ewald*volume);
+        eatom[i] *= qscale;
+      }
+    }
+
+    if (vflag_atom) {
+      for (i = 0; i < nlocal; i++)
+        for (j = 0; j < 6; j++) vatom[i][j] *= 0.5*q[i]*qscale;
+    }
+  }
+
   // sum energy across procs and add in volume-dependent term
 
   if (eflag_global) {
@@ -745,6 +778,10 @@ void PPPMGPU::poisson()
     work1[n++] *= scaleinv * greensfn[i];
   }
 
+  // extra FFTs for per-atom energy/virial
+
+  if (evflag_atom) poisson_peratom();
+
   // compute gradients of V(r) in each of 3 dims by transformimg -ik*V(k)
   // FFT leaves data in 3d brick decomposition
   // copy it into inner portion of vdx,vdy,vdz arrays
@@ -873,6 +910,12 @@ double PPPMGPU::memory_usage()
   bytes += nfft_both * sizeof(double);
   bytes += nfft_both*5 * sizeof(FFT_SCALAR);
   bytes += 2 * nbuf * sizeof(double);
+
+  if (peratom_allocate_flag) {
+    bytes += 7 * nbrick * sizeof(FFT_SCALAR);
+    bytes += 2 * nbuf_peratom * sizeof(FFT_SCALAR);
+  }
+
   return bytes + PPPM_GPU_API(bytes)();
 }
 

src/GPU/pppm_gpu.h

@@ -69,10 +69,10 @@ E: Cannot use order greater than 8 with pppm/gpu.
 
 Self-explanatory.
 
-E: Out of memory on GPGPU
+E: Insufficient memory on accelerator
 
-GPU memory is limited.  Reduce the size of the problem or increase the
+There is insufficient memory on one of the devices specified for the gpu
-number of GPUs.
+package
 
 E: Out of range atoms - cannot compute PPPM