Date: Tue, 17 May 2011 23:43:48 +0000
Subject: [PATCH 13/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6124
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/fix_adapt.html | 2 +-
doc/fix_adapt.txt | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/doc/fix_adapt.html b/doc/fix_adapt.html
index 837f898199..c26a9b2a39 100644
--- a/doc/fix_adapt.html
+++ b/doc/fix_adapt.html
@@ -54,7 +54,7 @@
fix 1 all adapt 1 pair soft a 1 1 v_prefactor
fix 1 all adapt 1 pair soft a 2* 3 v_prefactor
-fix 1 all adapt 1 pair lj/cut epsilon * * v_scale1 coul/cut pre 3 3 v_scale2 scale yes reset yes
+fix 1 all adapt 1 pair lj/cut epsilon * * v_scale1 coul/cut scale 3 3 v_scale2 scale yes reset yes
fix 1 all adapt 10 atom diameter v_size
Description:
diff --git a/doc/fix_adapt.txt b/doc/fix_adapt.txt
index f281c2cf01..3c6ecad3eb 100644
--- a/doc/fix_adapt.txt
+++ b/doc/fix_adapt.txt
@@ -41,7 +41,7 @@ keyword = {scale} or {reset} :l
fix 1 all adapt 1 pair soft a 1 1 v_prefactor
fix 1 all adapt 1 pair soft a 2* 3 v_prefactor
-fix 1 all adapt 1 pair lj/cut epsilon * * v_scale1 coul/cut pre 3 3 v_scale2 scale yes reset yes
+fix 1 all adapt 1 pair lj/cut epsilon * * v_scale1 coul/cut scale 3 3 v_scale2 scale yes reset yes
fix 1 all adapt 10 atom diameter v_size :pre
[Description:]
From 1e9c70d70c41a8800e78b98cd5e2b8621eaac512 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Tue, 17 May 2011 23:46:09 +0000
Subject: [PATCH 14/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6125
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index 30b4812a99..080b43a866 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "17 May 2011"
+#define LAMMPS_VERSION "18 May 2011"
From 38acfd9fadd35eaa6ab36164ba593a0dbb5c3cce Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Wed, 18 May 2011 14:25:33 +0000
Subject: [PATCH 15/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6128
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/group.cpp | 2 ++
1 file changed, 2 insertions(+)
diff --git a/src/group.cpp b/src/group.cpp
index b5803ccdb5..7ab12b35f3 100644
--- a/src/group.cpp
+++ b/src/group.cpp
@@ -142,6 +142,7 @@ void Group::assign(int narg, char **arg)
// style = region
// add to group if atom is in region
+ // init all regions via domain->init() to insure region can perform match()
if (strcmp(arg[1],"region") == 0) {
@@ -149,6 +150,7 @@ void Group::assign(int narg, char **arg)
int iregion = domain->find_region(arg[2]);
if (iregion == -1) error->all("Group region ID does not exist");
+ domain->init();
for (i = 0; i < nlocal; i++)
if (domain->regions[iregion]->match(x[i][0],x[i][1],x[i][2]))
From 14a348470c1975aa22f1a6cd646721222429923c Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Wed, 18 May 2011 20:03:24 +0000
Subject: [PATCH 16/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6136
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
lib/gpu/Makefile.linux | 4 ++
lib/gpu/README | 52 ++++++++++++++-----------
lib/gpu/cmm_cut_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/cmmc_long_gpu_kernel.cu | 64 +++++++++++++++----------------
lib/gpu/crml_gpu_kernel.cu | 68 ++++++++++++++++-----------------
lib/gpu/gb_gpu_kernel_nbor.cu | 20 +++++-----
lib/gpu/geryon/ucl_nv_kernel.h | 8 ++++
lib/gpu/lj96_cut_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/lj_cut_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/lj_expand_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/ljc_cut_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/ljcl_cut_gpu_kernel.cu | 64 +++++++++++++++----------------
lib/gpu/morse_gpu_kernel.cu | 48 +++++++++++------------
lib/gpu/pair_gpu_atom_kernel.cu | 14 +++----
lib/gpu/pair_gpu_device.cpp | 3 +-
lib/gpu/pair_gpu_device.h | 3 ++
lib/gpu/pppm_gpu_kernel.cu | 48 +++++++++++++----------
lib/gpu/pppm_gpu_memory.cpp | 6 ++-
18 files changed, 337 insertions(+), 305 deletions(-)
diff --git a/lib/gpu/Makefile.linux b/lib/gpu/Makefile.linux
index d69a00a817..1777187010 100644
--- a/lib/gpu/Makefile.linux
+++ b/lib/gpu/Makefile.linux
@@ -20,7 +20,11 @@
CUDA_HOME = /usr/local/cuda
NVCC = nvcc
+# newer CUDA
CUDA_ARCH = -arch=sm_13
+# older CUDA
+#CUDA_ARCH = -arch=sm_10 -DCUDA_PRE_THREE
+
CUDA_PRECISION = -D_SINGLE_SINGLE
CUDA_INCLUDE = -I$(CUDA_HOME)/include
CUDA_LIB = -L$(CUDA_HOME)/lib64
diff --git a/lib/gpu/README b/lib/gpu/README
index a60d43064a..73a51fc391 100644
--- a/lib/gpu/README
+++ b/lib/gpu/README
@@ -33,13 +33,17 @@ NOTE: Installation of the CUDA SDK is not required.
Current pair styles supporting GPU acceleration:
- 1. lj/cut/gpu
- 2. lj/cut/coul/cut/gpu
- 3. lj/cut/coul/long/gpu
- 4. lj96/cut/gpu
- 5. gayberne/gpu
- 6. cmm/cg/gpu
- 7. cmm/cg/coul/long/gpu
+ 1. lj/cut
+ 2. lj96/cut
+ 3. lj/expand
+ 4. lj/cut/coul/cut
+ 5. lj/cut/coul/long
+ 6. lj/charmm/coul/long
+ 7. morse
+ 8. cg/cmm
+ 9. cg/cmm/coul/long
+ 10. gayberne
+ 11. pppm
MULTIPLE LAMMPS PROCESSES
@@ -52,12 +56,12 @@ LAMMPS user manual for details on running with GPU acceleration.
BUILDING AND PRECISION MODES
-To build, edit the CUDA_ARCH, CUDA_PRECISION, CUDA_HOME, NVCC, CUDA_INCLUD,
-CUDA_LIB and CUDA_OPTS variables in one of the Makefiles. CUDA_ARCH should
-be set based on the compute capability of your GPU. This can be verified by
-running the nvc_get_devices executable after the build is complete.
-Additionally, the GPU package must be installed and compiled for LAMMPS.
-This may require editing the gpu_SYSPATH variable in the LAMMPS makefile.
+To build, edit the CUDA_ARCH, CUDA_PRECISION, CUDA_HOME variables in one of
+the Makefiles. CUDA_ARCH should be set based on the compute capability of
+your GPU. This can be verified by running the nvc_get_devices executable after
+the build is complete. Additionally, the GPU package must be installed and
+compiled for LAMMPS. This may require editing the gpu_SYSPATH variable in the
+LAMMPS makefile.
Please note that the GPU library accesses the CUDA driver library directly,
so it needs to be linked not only to the CUDA runtime library (libcudart.so)
@@ -74,6 +78,10 @@ the CUDA_PRECISION variable:
CUDA_PREC = -D_DOUBLE_DOUBLE # Double precision for all calculations
CUDA_PREC = -D_SINGLE_DOUBLE # Accumulation of forces, etc. in double
+NOTE: PPPM acceleration can only be run on GPUs with compute capability>=1.1.
+ You will get the error "GPU library not compiled for this accelerator."
+ when attempting to run PPPM on a GPU with compute capability 1.0.
+
NOTE: Double precision is only supported on certain GPUs (with
compute capability>=1.3).
@@ -83,15 +91,17 @@ NOTE: For Tesla and other graphics cards with compute capability>=1.3,
NOTE: For Fermi, make sure that -arch=sm_20 is set on the CUDA_ARCH line.
NOTE: The gayberne/gpu pair style will only be installed if the ASPHERE
- package has been installed before installing the GPU package in LAMMPS.
+ package has been installed.
NOTE: The cg/cmm/gpu and cg/cmm/coul/long/gpu pair styles will only be
- installed if the USER-CG-CMM package has been installed before
- installing the GPU package in LAMMPS.
+ installed if the USER-CG-CMM package has been installed.
-NOTE: The lj/cut/coul/long/gpu and cg/cmm/coul/long/gpu style will only be
- installed if the KSPACE package has been installed before installing
- the GPU package in LAMMPS.
+NOTE: The lj/cut/coul/long/gpu, cg/cmm/coul/long/gpu, pppm/gpu/single, and
+ pppm/gpu/double styles will only be installed if the KSPACE package has
+ been installed.
+
+NOTE: The lj/charmm/coul/long will only be installed if the MOLECULE package
+ has been installed.
EXAMPLE BUILD PROCESS
@@ -105,7 +115,3 @@ make yes-asphere
make yes-kspace
make yes-gpu
make linux
-
-------------------------------------------------------------------------
-Last merge with gpulammps: r561 on 2010-11-12
-------------------------------------------------------------------------
diff --git a/lib/gpu/cmm_cut_gpu_kernel.cu b/lib/gpu/cmm_cut_gpu_kernel.cu
index 08cc31ed7f..f99e7f06ac 100644
--- a/lib/gpu/cmm_cut_gpu_kernel.cu
+++ b/lib/gpu/cmm_cut_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef CMM_GPU_KERNEL
#define CMM_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -75,6 +51,30 @@ __inline float4 fetch_pos(const int& i, const float4 *pos)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/cmmc_long_gpu_kernel.cu b/lib/gpu/cmmc_long_gpu_kernel.cu
index 5153cb5016..a47a9267a1 100644
--- a/lib/gpu/cmmc_long_gpu_kernel.cu
+++ b/lib/gpu/cmmc_long_gpu_kernel.cu
@@ -18,38 +18,6 @@
#ifndef CMML_GPU_KERNEL
#define CMML_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
-#define EWALD_F (numtyp)1.12837917
-#define EWALD_P (numtyp)0.3275911
-#define A1 (numtyp)0.254829592
-#define A2 (numtyp)-0.284496736
-#define A3 (numtyp)1.421413741
-#define A4 (numtyp)-1.453152027
-#define A5 (numtyp)1.061405429
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -93,6 +61,38 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
+#define EWALD_F (numtyp)1.12837917
+#define EWALD_P (numtyp)0.3275911
+#define A1 (numtyp)0.254829592
+#define A2 (numtyp)-0.284496736
+#define A3 (numtyp)1.421413741
+#define A4 (numtyp)-1.453152027
+#define A5 (numtyp)1.061405429
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/crml_gpu_kernel.cu b/lib/gpu/crml_gpu_kernel.cu
index 63ce924581..dfdc7af3cd 100644
--- a/lib/gpu/crml_gpu_kernel.cu
+++ b/lib/gpu/crml_gpu_kernel.cu
@@ -18,40 +18,6 @@
#ifndef CRML_GPU_KERNEL
#define CRML_GPU_KERNEL
-#define MAX_BIO_SHARED_TYPES 128
-
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
-#define EWALD_F (numtyp)1.12837917
-#define EWALD_P (numtyp)0.3275911
-#define A1 (numtyp)0.254829592
-#define A2 (numtyp)-0.284496736
-#define A3 (numtyp)1.421413741
-#define A4 (numtyp)-1.453152027
-#define A5 (numtyp)1.061405429
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -94,6 +60,40 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+#define MAX_BIO_SHARED_TYPES 128
+
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
+#define EWALD_F (numtyp)1.12837917
+#define EWALD_P (numtyp)0.3275911
+#define A1 (numtyp)0.254829592
+#define A2 (numtyp)-0.284496736
+#define A3 (numtyp)1.421413741
+#define A4 (numtyp)-1.453152027
+#define A5 (numtyp)1.061405429
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/gb_gpu_kernel_nbor.cu b/lib/gpu/gb_gpu_kernel_nbor.cu
index 1b1d81fa42..b35b5b6998 100644
--- a/lib/gpu/gb_gpu_kernel_nbor.cu
+++ b/lib/gpu/gb_gpu_kernel_nbor.cu
@@ -18,16 +18,6 @@
#ifndef PAIR_GPU_KERNEL_H
#define PAIR_GPU_KERNEL_H
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#else
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -44,6 +34,16 @@
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#else
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#endif
+
// ---------------------------------------------------------------------------
// Unpack neighbors from dev_ij array into dev_nbor matrix for coalesced access
// -- Only unpack neighbors matching the specified inclusive range of forms
diff --git a/lib/gpu/geryon/ucl_nv_kernel.h b/lib/gpu/geryon/ucl_nv_kernel.h
index 5c45dc3a87..65a51b5f04 100644
--- a/lib/gpu/geryon/ucl_nv_kernel.h
+++ b/lib/gpu/geryon/ucl_nv_kernel.h
@@ -33,6 +33,14 @@
#define MEM_THREADS 32
#endif
+#ifdef CUDA_PRE_THREE
+struct __builtin_align__(16) _double4
+{
+ double x, y, z, w;
+};
+typedef struct _double4 double4;
+#endif
+
#define GLOBAL_ID_X threadIdx.x+mul24(blockIdx.x,blockDim.x)
#define GLOBAL_ID_Y threadIdx.y+mul24(blockIdx.y,blockDim.y)
#define GLOBAL_SIZE_X mul24(gridDim.x,blockDim.x);
diff --git a/lib/gpu/lj96_cut_gpu_kernel.cu b/lib/gpu/lj96_cut_gpu_kernel.cu
index 3fc6a2f308..1de9a8a7bf 100644
--- a/lib/gpu/lj96_cut_gpu_kernel.cu
+++ b/lib/gpu/lj96_cut_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef LJ96_GPU_KERNEL
#define LJ96_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -75,6 +51,30 @@ __inline float4 fetch_pos(const int& i, const float4 *pos)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/lj_cut_gpu_kernel.cu b/lib/gpu/lj_cut_gpu_kernel.cu
index 75f36446f7..9ef698cd09 100644
--- a/lib/gpu/lj_cut_gpu_kernel.cu
+++ b/lib/gpu/lj_cut_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef LJ_GPU_KERNEL
#define LJ_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -75,6 +51,30 @@ __inline float4 fetch_pos(const int& i, const float4 *pos)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/lj_expand_gpu_kernel.cu b/lib/gpu/lj_expand_gpu_kernel.cu
index 2d09b4d941..26fbefacf8 100644
--- a/lib/gpu/lj_expand_gpu_kernel.cu
+++ b/lib/gpu/lj_expand_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef LJE_GPU_KERNEL
#define LJE_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -75,6 +51,30 @@ __inline float4 fetch_pos(const int& i, const float4 *pos)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/ljc_cut_gpu_kernel.cu b/lib/gpu/ljc_cut_gpu_kernel.cu
index 44a607588a..ad1e530712 100644
--- a/lib/gpu/ljc_cut_gpu_kernel.cu
+++ b/lib/gpu/ljc_cut_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef LJC_GPU_KERNEL
#define LJC_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -85,6 +61,30 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/ljcl_cut_gpu_kernel.cu b/lib/gpu/ljcl_cut_gpu_kernel.cu
index 7be7a86114..ddde1dec32 100644
--- a/lib/gpu/ljcl_cut_gpu_kernel.cu
+++ b/lib/gpu/ljcl_cut_gpu_kernel.cu
@@ -18,38 +18,6 @@
#ifndef LJCL_GPU_KERNEL
#define LJCL_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
-#define EWALD_F (numtyp)1.12837917
-#define EWALD_P (numtyp)0.3275911
-#define A1 (numtyp)0.254829592
-#define A2 (numtyp)-0.284496736
-#define A3 (numtyp)1.421413741
-#define A4 (numtyp)-1.453152027
-#define A5 (numtyp)1.061405429
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -93,6 +61,38 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
+#define EWALD_F (numtyp)1.12837917
+#define EWALD_P (numtyp)0.3275911
+#define A1 (numtyp)0.254829592
+#define A2 (numtyp)-0.284496736
+#define A3 (numtyp)1.421413741
+#define A4 (numtyp)-1.453152027
+#define A5 (numtyp)1.061405429
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/morse_gpu_kernel.cu b/lib/gpu/morse_gpu_kernel.cu
index 0a89aae070..8832f58c64 100644
--- a/lib/gpu/morse_gpu_kernel.cu
+++ b/lib/gpu/morse_gpu_kernel.cu
@@ -18,30 +18,6 @@
#ifndef MORSE_GPU_KERNEL
#define MORSE_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp2 double2
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp2 float2
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "nv_kernel_def.h"
@@ -75,6 +51,30 @@ __inline float4 fetch_pos(const int& i, const float4 *pos)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp2 double2
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp2 float2
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
#define SBBITS 30
#define NEIGHMASK 0x3FFFFFFF
__inline int sbmask(int j) { return j >> SBBITS & 3; }
diff --git a/lib/gpu/pair_gpu_atom_kernel.cu b/lib/gpu/pair_gpu_atom_kernel.cu
index 2d1a6ba85f..ab79ac6e9c 100644
--- a/lib/gpu/pair_gpu_atom_kernel.cu
+++ b/lib/gpu/pair_gpu_atom_kernel.cu
@@ -15,6 +15,13 @@
Contributing authors: Mike Brown (ORNL), brownw@ornl.gov
------------------------------------------------------------------------- */
+#ifdef NV_KERNEL
+#include "geryon/ucl_nv_kernel.h"
+#else
+#pragma OPENCL EXTENSION cl_khr_fp64: enable
+#define GLOBAL_ID_X get_global_id(0)
+#endif
+
#ifdef _DOUBLE_DOUBLE
#define numtyp double
#define numtyp4 double4
@@ -23,13 +30,6 @@
#define numtyp4 float4
#endif
-#ifdef NV_KERNEL
-#include "geryon/ucl_nv_kernel.h"
-#else
-#pragma OPENCL EXTENSION cl_khr_fp64: enable
-#define GLOBAL_ID_X get_global_id(0)
-#endif
-
__kernel void kernel_cast_x(__global numtyp4 *x_type, __global double *x,
__global int *type, const int nall) {
int ii=GLOBAL_ID_X;
diff --git a/lib/gpu/pair_gpu_device.cpp b/lib/gpu/pair_gpu_device.cpp
index d5906b10e5..165d202832 100644
--- a/lib/gpu/pair_gpu_device.cpp
+++ b/lib/gpu/pair_gpu_device.cpp
@@ -549,8 +549,9 @@ int PairGPUDeviceT::compile_kernels() {
k_info.run(&d_gpu_lib_data.begin());
ucl_copy(h_gpu_lib_data,d_gpu_lib_data,false);
+ _ptx_arch=static_cast(h_gpu_lib_data[0])/100.0;
#ifndef USE_OPENCL
- if (static_cast(h_gpu_lib_data[0])/100.0>gpu->arch())
+ if (_ptx_arch>gpu->arch())
return -4;
#endif
diff --git a/lib/gpu/pair_gpu_device.h b/lib/gpu/pair_gpu_device.h
index 1e7e15e6a8..52b35cfcf2 100644
--- a/lib/gpu/pair_gpu_device.h
+++ b/lib/gpu/pair_gpu_device.h
@@ -226,6 +226,8 @@ class PairGPUDevice {
inline int block_bio_pair() const { return _block_bio_pair; }
/// Return the maximum number of atom types for shared mem with "bio" styles
inline int max_bio_shared_types() const { return _max_bio_shared_types; }
+ /// Architecture gpu code compiled for (returns 0 for OpenCL)
+ inline double ptx_arch() const { return _ptx_arch; }
// -------------------- SHARED DEVICE ROUTINES --------------------
// Perform asynchronous zero of integer array
@@ -281,6 +283,7 @@ class PairGPUDevice {
int _gpu_mode, _first_device, _last_device, _nthreads;
double _particle_split;
double _cpu_full;
+ double _ptx_arch;
int _num_mem_threads, _warp_size, _threads_per_atom, _threads_per_charge;
int _pppm_max_spline, _pppm_block;
diff --git a/lib/gpu/pppm_gpu_kernel.cu b/lib/gpu/pppm_gpu_kernel.cu
index c04e784de8..fe1862d051 100644
--- a/lib/gpu/pppm_gpu_kernel.cu
+++ b/lib/gpu/pppm_gpu_kernel.cu
@@ -18,27 +18,6 @@
#ifndef PPPM_GPU_KERNEL
#define PPPM_GPU_KERNEL
-#ifdef _DOUBLE_DOUBLE
-#define numtyp double
-#define numtyp4 double4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifdef _SINGLE_DOUBLE
-#define numtyp float
-#define numtyp4 float4
-#define acctyp double
-#define acctyp4 double4
-#endif
-
-#ifndef numtyp
-#define numtyp float
-#define numtyp4 float4
-#define acctyp float
-#define acctyp4 float4
-#endif
-
#ifdef NV_KERNEL
#include "geryon/ucl_nv_kernel.h"
@@ -67,6 +46,12 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+// Allow PPPM to compile without atomics for NVIDIA 1.0 cards, error
+// generated at runtime with use of pppm/gpu
+#if (__CUDA_ARCH__ < 110)
+#define atom_add(x,y) 0
+#endif
+
#else
#pragma OPENCL EXTENSION cl_khr_fp64: enable
@@ -85,6 +70,27 @@ __inline float fetch_q(const int& i, const float *q)
#endif
+#ifdef _DOUBLE_DOUBLE
+#define numtyp double
+#define numtyp4 double4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifdef _SINGLE_DOUBLE
+#define numtyp float
+#define numtyp4 float4
+#define acctyp double
+#define acctyp4 double4
+#endif
+
+#ifndef numtyp
+#define numtyp float
+#define numtyp4 float4
+#define acctyp float
+#define acctyp4 float4
+#endif
+
// Maximum order for spline
#define PPPM_MAX_SPLINE 8
// Thread block size for PPPM kernels
diff --git a/lib/gpu/pppm_gpu_memory.cpp b/lib/gpu/pppm_gpu_memory.cpp
index 521b3b1e46..2f7b35d051 100644
--- a/lib/gpu/pppm_gpu_memory.cpp
+++ b/lib/gpu/pppm_gpu_memory.cpp
@@ -66,7 +66,11 @@ grdtyp * PPPMGPUMemoryT::init(const int nlocal, const int nall, FILE *_screen,
flag=-5;
return 0;
}
-
+ if (device->ptx_arch()>0.0 && device->ptx_arch()<1.1) {
+ flag=-4;
+ return 0;
+ }
+
ucl_device=device->gpu;
atom=&device->atom;
From 3ac6369135104623c62067422a8d453ed076106d Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Thu, 19 May 2011 17:33:02 +0000
Subject: [PATCH 17/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6141
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/OPT/pair_eam_opt.h | 4 +++-
src/OPT/pair_lj_charmm_coul_long_opt.h | 11 ++++++-----
src/OPT/pair_lj_cut_opt.h | 9 +++++----
src/OPT/pair_morse_opt.h | 9 +++++----
4 files changed, 19 insertions(+), 14 deletions(-)
diff --git a/src/OPT/pair_eam_opt.h b/src/OPT/pair_eam_opt.h
index ddc074b316..d2318981f3 100644
--- a/src/OPT/pair_eam_opt.h
+++ b/src/OPT/pair_eam_opt.h
@@ -193,6 +193,7 @@ void PairEAMOpt::eval()
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
+ j &= NEIGHMASK;
double delx = xtmp - xx[j].x;
double dely = ytmp - xx[j].y;
@@ -269,7 +270,8 @@ void PairEAMOpt::eval()
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
-
+ j &= NEIGHMASK;
+
double delx = xtmp - xx[j].x;
double dely = ytmp - xx[j].y;
double delz = ztmp - xx[j].z;
diff --git a/src/OPT/pair_lj_charmm_coul_long_opt.h b/src/OPT/pair_lj_charmm_coul_long_opt.h
index 583f2d57e0..9cf5f8477b 100644
--- a/src/OPT/pair_lj_charmm_coul_long_opt.h
+++ b/src/OPT/pair_lj_charmm_coul_long_opt.h
@@ -66,7 +66,7 @@ void PairLJCharmmCoulLongOpt::eval()
double _pad[2];
} fast_alpha_t;
- int i,j,ii,jj,inum,jnum,itype,jtype,itable;
+ int i,j,ii,jj,inum,jnum,itype,jtype,itable,sbindex;
double fraction,table;
double r,r2inv,r6inv,forcecoul,forcelj,factor_coul,factor_lj;
double grij,expm2,prefactor,t,erfc;
@@ -132,8 +132,9 @@ void PairLJCharmmCoulLongOpt::eval()
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
-
- if (j <= NEIGHMASK) {
+ sbindex = sbmask(j);
+
+ if (sbindex == 0) {
double delx = xtmp - xx[j].x;
double dely = ytmp - xx[j].y;
double delz = ztmp - xx[j].z;
@@ -219,8 +220,8 @@ void PairLJCharmmCoulLongOpt::eval()
}
} else {
- factor_lj = special_lj[sbmask(j)];
- factor_coul = special_coul[sbmask(j)];
+ factor_lj = special_lj[sbindex];
+ factor_coul = special_coul[sbindex];
j &= NEIGHMASK;
double delx = xtmp - xx[j].x;
diff --git a/src/OPT/pair_lj_cut_opt.h b/src/OPT/pair_lj_cut_opt.h
index dce90b0249..ae7ac7d7af 100644
--- a/src/OPT/pair_lj_cut_opt.h
+++ b/src/OPT/pair_lj_cut_opt.h
@@ -54,7 +54,8 @@ void PairLJCutOpt::eval()
double _pad[2];
} fast_alpha_t;
- int i,j,ii,jj,inum,jnum,itype,jtype;
+ int i,j,ii,jj,inum,jnum,itype,jtype,sbindex;
+ double factor_lj;
double evdwl = 0.0;
double** __restrict__ x = atom->x;
@@ -106,9 +107,9 @@ void PairLJCutOpt::eval()
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
- double factor_lj;
+ sbindex = sbmask(j);
- if (j <= NEIGHMASK) {
+ if (sbindex == 0) {
double delx = xtmp - xx[j].x;
double dely = ytmp - xx[j].y;
double delz = ztmp - xx[j].z;
@@ -141,7 +142,7 @@ void PairLJCutOpt::eval()
}
} else {
- factor_lj = special_lj[sbmask(j)];
+ factor_lj = special_lj[sbindex];
j &= NEIGHMASK;
double delx = xtmp - xx[j].x;
diff --git a/src/OPT/pair_morse_opt.h b/src/OPT/pair_morse_opt.h
index 0e92408e1d..f7e9710433 100644
--- a/src/OPT/pair_morse_opt.h
+++ b/src/OPT/pair_morse_opt.h
@@ -55,7 +55,8 @@ void PairMorseOpt::eval()
double _pad[2];
} fast_alpha_t;
- int i,j,ii,jj,inum,jnum,itype,jtype;
+ int i,j,ii,jj,inum,jnum,itype,jtype,sbindex;
+ double factor_lj;
double evdwl = 0.0;
double** __restrict__ x = atom->x;
@@ -107,9 +108,9 @@ void PairMorseOpt::eval()
for (jj = 0; jj < jnum; jj++) {
j = jlist[jj];
- double factor_lj;
+ sbindex = sbmask(j);
- if (j <= NEIGHMASK) {
+ if (sbindex == 0) {
double delx = xtmp - xx[j].x;
double dely = ytmp - xx[j].y;
double delz = ztmp - xx[j].z;
@@ -140,7 +141,7 @@ void PairMorseOpt::eval()
}
} else {
- factor_lj = special_lj[sbmask(j)];
+ factor_lj = special_lj[sbindex];
j &= NEIGHMASK;
double delx = xtmp - xx[j].x;
From ac93eef58feb984b5a1deae4d2b6a363016820cd Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Thu, 19 May 2011 17:35:14 +0000
Subject: [PATCH 18/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6142
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index 080b43a866..110bee8220 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "18 May 2011"
+#define LAMMPS_VERSION "19 May 2011"
From ff9b68ba359ada2ab5b18222179d7cb190fc35ae Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:19:48 +0000
Subject: [PATCH 19/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6146
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/MAKE/Makefile.g++ | 19 ++++++++++++++-----
src/Makefile | 2 +-
2 files changed, 15 insertions(+), 6 deletions(-)
diff --git a/src/MAKE/Makefile.g++ b/src/MAKE/Makefile.g++
index 147ac4f388..25144f903a 100755
--- a/src/MAKE/Makefile.g++
+++ b/src/MAKE/Makefile.g++
@@ -38,7 +38,7 @@ MPI_LIB = -lmpich -lpthread
# PATH = path for FFT library
# LIB = name of FFT library
-FFT_INC = -DFFT_FFTW
+FFT_INC = -DFFT_FFTW
FFT_PATH =
FFT_LIB = -lfftw
@@ -51,21 +51,30 @@ JPG_INC =
JPG_PATH =
JPG_LIB =
-# additional system libraries needed by LAMMPS package libraries
+# additional system settings needed by LAMMPS package libraries
# these settings are IGNORED if the corresponding LAMMPS package
# (e.g. gpu, meam) is NOT included in the LAMMPS build
-# SYSLIB = names of libraries
-# SYSPATH = paths of libraries
+# SYSINC = settings to compile with
+# SYSLIB = libraries to link with
+# SYSPATH = paths to libraries
+
+gpu_SYSINC =
+meam_SYSINC =
+reax_SYSINC =
+user-atc_SYSINC =
+user-cuda_SYSINC = -I/usr/local/cuda/include -DCUDA -DCUDA_ARCH=20 -DFFT_CUFFT
gpu_SYSLIB = -lcudart -lcuda
meam_SYSLIB = -lgfortran
reax_SYSLIB = -lgfortran
user-atc_SYSLIB = -lblas -llapack
+user-cuda_SYSLIB = -lcudart -lcuda
gpu_SYSPATH = -L/usr/local/cuda/lib64
meam_SYSPATH =
reax_SYSPATH =
user-atc_SYSPATH =
+user-cuda_SYSPATH =
# ---------------------------------------------------------------------
# build rules and dependencies
@@ -73,7 +82,7 @@ user-atc_SYSPATH =
include Makefile.package
-EXTRA_INC = $(LMP_INC) $(PKG_INC) $(MPI_INC) $(FFT_INC) $(JPG_INC)
+EXTRA_INC = $(LMP_INC) $(PKG_INC) $(MPI_INC) $(FFT_INC) $(JPG_INC) $(PKG_SYSINC)
EXTRA_PATH = $(PKG_PATH) $(MPI_PATH) $(FFT_PATH) $(JPG_PATH) $(PKG_SYSPATH)
EXTRA_LIB = $(PKG_LIB) $(MPI_LIB) $(FFT_LIB) $(JPG_LIB) $(PKG_SYSLIB)
diff --git a/src/Makefile b/src/Makefile
index 7473b0e112..621a446289 100755
--- a/src/Makefile
+++ b/src/Makefile
@@ -17,7 +17,7 @@ PACKAGE = asphere class2 colloid dipole dsmc gpu granular \
kspace manybody meam molecule opt peri poems reax replica \
shock srd xtc
-PACKUSER = user-ackland user-atc user-cd-eam user-cg-cmm user-eff \
+PACKUSER = user-ackland user-atc user-cd-eam user-cg-cmm user-cuda user-eff \
user-ewaldn user-imd user-reaxc user-smd
PACKALL = $(PACKAGE) $(PACKUSER)
From ebcb5d660b8a991588ffb559790b24acb2f6dc3c Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:20:03 +0000
Subject: [PATCH 20/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6147
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/Makefile.package.empty | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/src/Makefile.package.empty b/src/Makefile.package.empty
index fbee509ca6..ef54f586dc 100644
--- a/src/Makefile.package.empty
+++ b/src/Makefile.package.empty
@@ -5,5 +5,6 @@ PKG_INC =
PKG_PATH =
PKG_LIB =
-PKG_SYSPATH =
+PKG_SYSINC =
PKG_SYSLIB =
+PKG_SYSPATH =
From e4a549ee7bda6ae20bb0f0620ab23c1c66433d32 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:20:27 +0000
Subject: [PATCH 21/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6148
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/GPU/Install.sh | 9 +++++++--
src/MEAM/Install.sh | 3 ++-
src/REAX/Install.sh | 3 ++-
3 files changed, 11 insertions(+), 4 deletions(-)
diff --git a/src/GPU/Install.sh b/src/GPU/Install.sh
index d338b516a5..6367a6cea5 100644
--- a/src/GPU/Install.sh
+++ b/src/GPU/Install.sh
@@ -10,8 +10,9 @@ if (test $1 = 1) then
sed -i -e 's/[^ \t]*gpu_[^ \t]*) //' ../Makefile.package
sed -i -e 's|^PKG_PATH =[ \t]*|&-L../../lib/gpu |' ../Makefile.package
sed -i -e 's|^PKG_LIB =[ \t]*|&-lgpu |' ../Makefile.package
- sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(gpu_SYSPATH) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSINC =[ \t]*|&$(gpu_SYSINC) |' ../Makefile.package
sed -i -e 's|^PKG_SYSLIB =[ \t]*|&$(gpu_SYSLIB) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(gpu_SYSPATH) |' ../Makefile.package
fi
if (test -e ../pppm.cpp) then
@@ -53,13 +54,15 @@ if (test $1 = 1) then
cp pair_lj96_cut_gpu.cpp ..
cp pair_lj_expand_gpu.cpp ..
cp pair_lj_cut_coul_cut_gpu.cpp ..
+
+ cp fix_gpu.cpp ..
+
cp pair_lj_cut_gpu.h ..
cp pair_morse_gpu.h ..
cp pair_lj96_cut_gpu.h ..
cp pair_lj_expand_gpu.h ..
cp pair_lj_cut_coul_cut_gpu.h ..
- cp fix_gpu.cpp ..
cp fix_gpu.h ..
cp gpu_extra.h ..
@@ -83,6 +86,7 @@ elif (test $1 = 0) then
rm ../pair_lj_charmm_coul_long_gpu.cpp
rm ../pair_cg_cmm_gpu.cpp
rm ../pair_cg_cmm_coul_long_gpu.cpp
+
rm ../fix_gpu.cpp
rm ../pppm_gpu.h
@@ -98,6 +102,7 @@ elif (test $1 = 0) then
rm ../pair_lj_charmm_coul_long_gpu.h
rm ../pair_cg_cmm_gpu.h
rm ../pair_cg_cmm_coul_long_gpu.h
+
rm ../fix_gpu.h
rm ../gpu_extra.h
diff --git a/src/MEAM/Install.sh b/src/MEAM/Install.sh
index e3c9d623f3..4ce114f3c4 100644
--- a/src/MEAM/Install.sh
+++ b/src/MEAM/Install.sh
@@ -9,8 +9,9 @@ if (test $1 = 1) then
sed -i -e 's|^PKG_INC =[ \t]*|&-I../../lib/meam |' ../Makefile.package
sed -i -e 's|^PKG_PATH =[ \t]*|&-L../../lib/meam |' ../Makefile.package
sed -i -e 's|^PKG_LIB =[ \t]*|&-lmeam |' ../Makefile.package
- sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(meam_SYSPATH) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSINC =[ \t]*|&$(meam_SYSINC) |' ../Makefile.package
sed -i -e 's|^PKG_SYSLIB =[ \t]*|&$(meam_SYSLIB) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(meam_SYSPATH) |' ../Makefile.package
fi
cp pair_meam.cpp ..
diff --git a/src/REAX/Install.sh b/src/REAX/Install.sh
index bd56f798db..abf9078568 100644
--- a/src/REAX/Install.sh
+++ b/src/REAX/Install.sh
@@ -9,8 +9,9 @@ if (test $1 = 1) then
sed -i -e 's|^PKG_INC =[ \t]*|&-I../../lib/reax |' ../Makefile.package
sed -i -e 's|^PKG_PATH =[ \t]*|&-L../../lib/reax |' ../Makefile.package
sed -i -e 's|^PKG_LIB =[ \t]*|&-lreax |' ../Makefile.package
- sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(reax_SYSPATH) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSINC =[ \t]*|&$(reax_SYSINC) |' ../Makefile.package
sed -i -e 's|^PKG_SYSLIB =[ \t]*|&$(reax_SYSLIB) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(reax_SYSPATH) |' ../Makefile.package
fi
cp pair_reax.cpp ..
From 40af1682eeb574b85f5939fbe0200ea6ab2a8b22 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:24:03 +0000
Subject: [PATCH 22/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6149
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/USER-ATC/Install.sh | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/src/USER-ATC/Install.sh b/src/USER-ATC/Install.sh
index e55e1c7d65..307b9cab00 100755
--- a/src/USER-ATC/Install.sh
+++ b/src/USER-ATC/Install.sh
@@ -9,8 +9,9 @@ if (test $1 = 1) then
sed -i -e 's|^PKG_INC =[ \t]*|&-I../../lib/atc |' ../Makefile.package
sed -i -e 's|^PKG_PATH =[ \t]*|&-L../../lib/atc |' ../Makefile.package
sed -i -e 's|^PKG_LIB =[ \t]*|&-latc |' ../Makefile.package
- sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(user-atc_SYSPATH) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSINC =[ \t]*|&$(user-atc_SYSINC) |' ../Makefile.package
sed -i -e 's|^PKG_SYSLIB =[ \t]*|&$(user-atc_SYSLIB) |' ../Makefile.package
+ sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(user-atc_SYSPATH) |' ../Makefile.package
fi
cp fix_atc.h ..
From 5b5d4db5ad395cfbc328bdec650d6146df3687a4 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:25:36 +0000
Subject: [PATCH 23/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6150
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/lammps.cpp | 64 ++++++++++++++++++++++++++++++++++++++------------
src/lammps.h | 4 ++++
2 files changed, 53 insertions(+), 15 deletions(-)
diff --git a/src/lammps.cpp b/src/lammps.cpp
index 15b32bac2f..9ccbcbde70 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -27,10 +27,13 @@
#include "modify.h"
#include "group.h"
#include "output.h"
+#include "accelerator.h"
#include "timer.h"
using namespace LAMMPS_NS;
+enum{NONE,OPT,GPU,USERCUDA};
+
/* ----------------------------------------------------------------------
start up LAMMPS
allocate fundamental classes (memory, error, universe, input)
@@ -54,14 +57,17 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
int inflag = 0;
int screenflag = 0;
int logflag = 0;
+ accelerator = NONE;
+ cuda = NULL;
+ asuffix = NULL;
+
int iarg = 1;
while (iarg < narg) {
if (strcmp(arg[iarg],"-partition") == 0 ||
strcmp(arg[iarg],"-p") == 0) {
universe->existflag = 1;
- if (iarg+2 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
iarg++;
while (iarg < narg && arg[iarg][0] != '-') {
universe->add_world(arg[iarg]);
@@ -69,33 +75,37 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
}
} else if (strcmp(arg[iarg],"-in") == 0 ||
strcmp(arg[iarg],"-i") == 0) {
- if (iarg+2 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
inflag = iarg + 1;
iarg += 2;
} else if (strcmp(arg[iarg],"-screen") == 0 ||
strcmp(arg[iarg],"-s") == 0) {
- if (iarg+2 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
screenflag = iarg + 1;
iarg += 2;
} else if (strcmp(arg[iarg],"-log") == 0 ||
strcmp(arg[iarg],"-l") == 0) {
- if (iarg+2 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
logflag = iarg + 1;
iarg += 2;
} else if (strcmp(arg[iarg],"-var") == 0 ||
strcmp(arg[iarg],"-v") == 0) {
- if (iarg+3 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+3 > narg) error->universe_all("Invalid command-line argument");
iarg += 2;
while (iarg < narg && arg[iarg][0] != '-') iarg++;
} else if (strcmp(arg[iarg],"-echo") == 0 ||
strcmp(arg[iarg],"-e") == 0) {
- if (iarg+2 > narg)
- error->universe_all("Invalid command-line argument");
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
iarg += 2;
+ } else if (strcmp(arg[iarg],"-accel") == 0 ||
+ strcmp(arg[iarg],"-a") == 0) {
+ if (iarg+2 > narg) error->universe_all("Invalid command-line argument");
+ if (strcmp(arg[iarg+1],"opt") == 0) accelerator = OPT;
+ else if (strcmp(arg[iarg+1],"gpu") == 0) accelerator = GPU;
+ else if (strcmp(arg[iarg+1],"cuda") == 0) accelerator = USERCUDA;
+ else error->universe_all("Invalid command-line argument");
+ asuffix = new char[8];
+ strcpy(asuffix,arg[iarg+1]);
} else error->universe_all("Invalid command-line argument");
}
@@ -265,6 +275,16 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
if (mpisize != sizeof(bigint))
error->all("MPI_LMP_BIGINT and bigint in lmptype.h are not compatible");
+ // check consistency of -a switch with installed packages
+ // for OPT and GPU, no problem if not installed
+ // for USER-CUDA, throw error if not installed
+
+ if (accelerator == USERCUDA) {
+ cuda = new Cuda(this);
+ if (!cuda->cuda_exists)
+ error->all("Command-line switch requires USER-CUDA package be installed");
+ }
+
// allocate input class now that MPI is fully setup
input = new Input(this,narg,arg);
@@ -285,6 +305,7 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
LAMMPS::~LAMMPS()
{
destroy();
+ if (accelerator == USERCUDA) delete cuda;
if (universe->nworlds == 1) {
if (logfile) fclose(logfile);
@@ -296,6 +317,8 @@ LAMMPS::~LAMMPS()
if (world != universe->uworld) MPI_Comm_free(&world);
+ delete [] asuffix;
+
delete input;
delete universe;
delete error;
@@ -305,17 +328,26 @@ LAMMPS::~LAMMPS()
/* ----------------------------------------------------------------------
allocate single instance of top-level classes
fundamental classes are allocated in constructor
+ some classes have accelerator variants
------------------------------------------------------------------------- */
void LAMMPS::create()
{
atom = new Atom(this);
neighbor = new Neighbor(this);
- comm = new Comm(this);
- domain = new Domain(this);
+
+ if (accelerator == USERCUDA) comm = new CommCuda(this);
+ else comm = new Comm(this);
+
+ if (accelerator == USERCUDA) domain = new DomainCuda(this);
+ else domain = new Domain(this);
+
group = new Group(this);
force = new Force(this); // must be after group, to create temperature
- modify = new Modify(this);
+
+ if (accelerator == USERCUDA) modify = new ModifyCuda(this);
+ else modify = new Modify(this);
+
output = new Output(this); // must be after group, so "all" exists
// must be after modify so can create Computes
update = new Update(this); // must be after output, force, neighbor
@@ -328,6 +360,8 @@ void LAMMPS::create()
void LAMMPS::init()
{
+ if (accelerator == USERCUDA) cuda->setDevice(this);
+
update->init();
force->init(); // pair must come after update due to minimizer
domain->init();
diff --git a/src/lammps.h b/src/lammps.h
index 1c3f81600e..095fa1054d 100644
--- a/src/lammps.h
+++ b/src/lammps.h
@@ -42,6 +42,10 @@ class LAMMPS {
FILE *screen; // screen output
FILE *logfile; // logfile
+ int accelerator; // accelerator flag
+ char *asuffix; // accelerator suffix
+ class Cuda *cuda; // CUDA accelerator class
+
LAMMPS(int, char **, MPI_Comm);
~LAMMPS();
void create();
From 8037991332fbd00eef3366f93ccd602da684b7b3 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:27:50 +0000
Subject: [PATCH 24/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6151
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/verlet.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/verlet.h b/src/verlet.h
index 211b5653f6..8b8c90b7a9 100644
--- a/src/verlet.h
+++ b/src/verlet.h
@@ -27,14 +27,14 @@ namespace LAMMPS_NS {
class Verlet : public Integrate {
public:
Verlet(class LAMMPS *, int, char **);
- ~Verlet() {}
+ virtual ~Verlet() {}
void init();
void setup();
void setup_minimal(int);
void run(int);
void cleanup();
- private:
+ protected:
int triclinic; // 0 if domain is orthog, 1 if triclinic
int torqueflag; // zero out arrays every step
int erforceflag;
From 3145254774787eb19aa85d02ad065b35a7085cf4 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:28:28 +0000
Subject: [PATCH 25/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6152
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/compute.cpp | 1 +
src/compute.h | 2 ++
2 files changed, 3 insertions(+)
diff --git a/src/compute.cpp b/src/compute.cpp
index 5c5398e5df..230c15919c 100644
--- a/src/compute.cpp
+++ b/src/compute.cpp
@@ -68,6 +68,7 @@ Compute::Compute(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)
timeflag = 0;
comm_forward = comm_reverse = 0;
+ cudable = 0;
invoked_scalar = invoked_vector = invoked_array = -1;
invoked_peratom = invoked_local = -1;
diff --git a/src/compute.h b/src/compute.h
index 9b929b345a..b0f9709df4 100644
--- a/src/compute.h
+++ b/src/compute.h
@@ -77,6 +77,8 @@ class Compute : protected Pointers {
int comm_forward; // size of forward communication (0 if none)
int comm_reverse; // size of reverse communication (0 if none)
+ int cudable; // 1 if compute is CUDA-enabled
+
Compute(class LAMMPS *, int, char **);
virtual ~Compute();
void modify_params(int, char **);
From f4d2ebfb545db98633306ec3efb49000be4a2095 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:40:22 +0000
Subject: [PATCH 26/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6153
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/modify.cpp | 67 +++++++++++++++++++++++++++++++++++++++++---------
src/modify.h | 8 +++---
2 files changed, 60 insertions(+), 15 deletions(-)
diff --git a/src/modify.cpp b/src/modify.cpp
index 329cd268a2..514341fc17 100644
--- a/src/modify.cpp
+++ b/src/modify.cpp
@@ -591,7 +591,7 @@ int Modify::min_reset_ref()
add a new fix or replace one with same ID
------------------------------------------------------------------------- */
-void Modify::add_fix(int narg, char **arg)
+void Modify::add_fix(int narg, char **arg, char *suffix)
{
if (domain->box_exist == 0)
error->all("Fix command before simulation box is defined");
@@ -636,17 +636,39 @@ void Modify::add_fix(int narg, char **arg)
}
}
- // create the Fix
+ // create the Fix, first with suffix appended
- if (0) return; // dummy line to enable else-if macro expansion
+ int success = 0;
+
+ if (suffix) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",arg[2],suffix);
+ success = 1;
+
+ if (0) return;
#define FIX_CLASS
#define FixStyle(key,Class) \
- else if (strcmp(arg[2],#key) == 0) fix[ifix] = new Class(lmp,narg,arg);
+ else if (strcmp(estyle,#key) == 0) fix[ifix] = new Class(lmp,narg,arg);
#include "style_fix.h"
+#undef FixStyle
#undef FIX_CLASS
- else error->all("Invalid fix style");
+ else success = 0;
+ }
+
+ if (!success) {
+ if (0) return;
+
+#define FIX_CLASS
+#define FixStyle(key,Class) \
+ else if (strcmp(arg[2],#key) == 0) fix[ifix] = new Class(lmp,narg,arg);
+#include "style_fix.h"
+#undef FixStyle
+#undef FIX_CLASS
+
+ else error->all("Invalid fix style");
+ }
// set fix mask values and increment nfix (if new)
@@ -740,7 +762,7 @@ int Modify::find_fix(const char *id)
add a new compute
------------------------------------------------------------------------- */
-void Modify::add_compute(int narg, char **arg)
+void Modify::add_compute(int narg, char **arg, char *suffix)
{
if (narg < 3) error->all("Illegal compute command");
@@ -758,18 +780,41 @@ void Modify::add_compute(int narg, char **arg)
memory->srealloc(compute,maxcompute*sizeof(Compute *),"modify:compute");
}
- // create the Compute
+ // create the Compute, first with suffix appended
- if (0) return; // dummy line to enable else-if macro expansion
+ int success = 0;
+
+ if (suffix) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",arg[2],suffix);
+ success = 1;
+
+ if (0) return;
#define COMPUTE_CLASS
#define ComputeStyle(key,Class) \
- else if (strcmp(arg[2],#key) == 0) \
- compute[ncompute] = new Class(lmp,narg,arg);
+ else if (strcmp(estyle,#key) == 0) \
+ compute[ncompute] = new Class(lmp,narg,arg);
#include "style_compute.h"
+#undef ComputeStyle
#undef COMPUTE_CLASS
- else error->all("Invalid compute style");
+ else success = 0;
+ }
+
+ if (!success) {
+ if (0) return;
+
+#define COMPUTE_CLASS
+#define ComputeStyle(key,Class) \
+ else if (strcmp(arg[2],#key) == 0) \
+ compute[ncompute] = new Class(lmp,narg,arg);
+#include "style_compute.h"
+#undef ComputeStyle
+#undef COMPUTE_CLASS
+
+ else error->all("Invalid compute style");
+ }
ncompute++;
}
diff --git a/src/modify.h b/src/modify.h
index 81d9401e69..fdd25b1045 100644
--- a/src/modify.h
+++ b/src/modify.h
@@ -40,7 +40,7 @@ class Modify : protected Pointers {
class Compute **compute;
Modify(class LAMMPS *);
- ~Modify();
+ virtual ~Modify();
void init();
void setup(int);
void setup_pre_exchange();
@@ -79,12 +79,12 @@ class Modify : protected Pointers {
double max_alpha(double *);
int min_dof();
- void add_fix(int, char **);
+ void add_fix(int, char **, char *suffix = NULL);
void modify_fix(int, char **);
void delete_fix(const char *);
int find_fix(const char *);
- void add_compute(int, char **);
+ void add_compute(int, char **, char *suffix = NULL);
void modify_compute(int, char **);
void delete_compute(char *);
int find_compute(char *);
@@ -98,7 +98,7 @@ class Modify : protected Pointers {
bigint memory_usage();
- private:
+ protected:
// lists of fixes to apply at different stages of timestep
From 93115ad750234a5cb28792675cdf24503283f6e7 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:44:23 +0000
Subject: [PATCH 27/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6154
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/force.cpp | 36 +++++++++++++++++++++++++++++-------
src/force.h | 4 ++--
2 files changed, 31 insertions(+), 9 deletions(-)
diff --git a/src/force.cpp b/src/force.cpp
index 8068717be5..867101dac4 100644
--- a/src/force.cpp
+++ b/src/force.cpp
@@ -117,32 +117,54 @@ void Force::init()
create a pair style, called from input script or restart file
------------------------------------------------------------------------- */
-void Force::create_pair(const char *style)
+void Force::create_pair(const char *style, char *suffix)
{
delete [] pair_style;
if (pair) delete pair;
- pair = new_pair(style);
+ pair = new_pair(style,suffix);
int n = strlen(style) + 1;
pair_style = new char[n];
strcpy(pair_style,style);
}
/* ----------------------------------------------------------------------
- generate a pair class
+ generate a pair class, first with suffix appended
------------------------------------------------------------------------- */
-Pair *Force::new_pair(const char *style)
+Pair *Force::new_pair(const char *style, char *suffix)
{
- if (strcmp(style,"none") == 0) return NULL;
+ int success = 0;
+
+ if (suffix) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",style,suffix);
+ success = 1;
+
+ if (0) return NULL;
#define PAIR_CLASS
#define PairStyle(key,Class) \
- else if (strcmp(style,#key) == 0) return new Class(lmp);
+ else if (strcmp(estyle,#key) == 0) return new Class(lmp);
+#include "style_pair.h"
+#undef PairStyle
+#undef PAIR_CLASS
+
+ else success = 0;
+ }
+
+ if (!success) {
+ if (strcmp(style,"none") == 0) return NULL;
+
+#define PAIR_CLASS
+#define PairStyle(key,Class) \
+ else if (strcmp(style,#key) == 0) return new Class(lmp);
#include "style_pair.h"
#undef PAIR_CLASS
- else error->all("Invalid pair style");
+ else error->all("Invalid pair style");
+ }
+
return NULL;
}
diff --git a/src/force.h b/src/force.h
index ee6343db98..d8b4a20136 100644
--- a/src/force.h
+++ b/src/force.h
@@ -64,8 +64,8 @@ class Force : protected Pointers {
~Force();
void init();
- void create_pair(const char *);
- class Pair *new_pair(const char *);
+ void create_pair(const char *, char *suffix = NULL);
+ class Pair *new_pair(const char *, char *suffix = NULL);
class Pair *pair_match(const char *, int);
void create_bond(const char *);
From 66e3ed7c5892bbd9be0dc23dac1a616a5acc75fc Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:46:27 +0000
Subject: [PATCH 28/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6155
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/pair.cpp | 5 +++++
1 file changed, 5 insertions(+)
diff --git a/src/pair.cpp b/src/pair.cpp
index b8331086f7..dde19c1d05 100644
--- a/src/pair.cpp
+++ b/src/pair.cpp
@@ -30,6 +30,7 @@
#include "comm.h"
#include "force.h"
#include "update.h"
+#include "accelerator.h"
#include "memory.h"
#include "error.h"
@@ -40,6 +41,7 @@ using namespace LAMMPS_NS;
enum{GEOMETRIC,ARITHMETIC,SIXTHPOWER};
enum{R,RSQ,BMP};
+enum{NOACCEL,OPT,GPU,USERCUDA}; // same as lammps.cpp
/* ---------------------------------------------------------------------- */
@@ -325,6 +327,9 @@ void Pair::ev_setup(int eflag, int vflag)
if (vflag_atom == 0) vflag_either = 0;
if (vflag_either == 0 && eflag_either == 0) evflag = 0;
} else vflag_fdotr = 0;
+
+ if (lmp->accelerator == USERCUDA)
+ lmp->cuda->evsetup_eatom_vatom(eflag_atom,vflag_atom);
}
/* ----------------------------------------------------------------------
From cef6d81fe6dd82d3579c0ac6a597b63e504956c5 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:50:14 +0000
Subject: [PATCH 29/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6156
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom_vec.h | 3 +++
1 file changed, 3 insertions(+)
diff --git a/src/atom_vec.h b/src/atom_vec.h
index a437d14a6b..35cae37bb6 100644
--- a/src/atom_vec.h
+++ b/src/atom_vec.h
@@ -38,6 +38,9 @@ class AtomVec : protected Pointers {
int size_data_bonus; // number of values in Bonus line
int xcol_data; // column (1-N) where x is in Atom line
+ int cudable; // 1 if atom style is CUDA-enabled
+ int *maxsend; // CUDA-specific variable
+
AtomVec(class LAMMPS *, int, char **);
virtual ~AtomVec() {}
virtual void init();
From ae07e039c577a00c5ca5c71c3cfec8ad9b096f0e Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:53:02 +0000
Subject: [PATCH 30/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6157
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/input.cpp | 24 ++++++++++++++++++++----
src/input.h | 3 ++-
2 files changed, 22 insertions(+), 5 deletions(-)
diff --git a/src/input.cpp b/src/input.cpp
index 0861c96cc2..de4dcb97e9 100644
--- a/src/input.cpp
+++ b/src/input.cpp
@@ -42,6 +42,7 @@
#include "neighbor.h"
#include "special.h"
#include "variable.h"
+#include "accelerator.h"
#include "error.h"
#include "memory.h"
@@ -417,6 +418,7 @@ int Input::execute_command()
else if (!strcmp(command,"shell")) shell();
else if (!strcmp(command,"variable")) variable_command();
+ else if (!strcmp(command,"accelerator")) accelerator();
else if (!strcmp(command,"angle_coeff")) angle_coeff();
else if (!strcmp(command,"angle_style")) angle_style();
else if (!strcmp(command,"atom_modify")) atom_modify();
@@ -801,6 +803,20 @@ void Input::variable_command()
one function for each LAMMPS-specific input script command
------------------------------------------------------------------------- */
+void Input::accelerator()
+{
+ if (domain->box_exist)
+ error->all("Accelerator command after simulation box is defined");
+ if (narg < 1) error->all("Illegal accelerator command");
+ if (strcmp(lmp->asuffix,arg[0]) != 0)
+ error->all("Accelerator command requires matching command-line -a switch");
+
+ if (strcmp(arg[0],"cuda") == 0) lmp->cuda->accelerator(narg-1,&arg[1]);
+ else error->all("Illegal accelerator command");
+}
+
+/* ---------------------------------------------------------------------- */
+
void Input::angle_coeff()
{
if (domain->box_exist == 0)
@@ -837,7 +853,7 @@ void Input::atom_style()
if (narg < 1) error->all("Illegal atom_style command");
if (domain->box_exist)
error->all("Atom_style command after simulation box is defined");
- atom->create_avec(arg[0],narg-1,&arg[1]);
+ atom->create_avec(arg[0],narg-1,&arg[1],lmp->asuffix);
}
/* ---------------------------------------------------------------------- */
@@ -884,7 +900,7 @@ void Input::communicate()
void Input::compute()
{
- modify->add_compute(narg,arg);
+ modify->add_compute(narg,arg,lmp->asuffix);
}
/* ---------------------------------------------------------------------- */
@@ -962,7 +978,7 @@ void Input::dump_modify()
void Input::fix()
{
- modify->add_fix(narg,arg);
+ modify->add_fix(narg,arg,lmp->asuffix);
}
/* ---------------------------------------------------------------------- */
@@ -1132,7 +1148,7 @@ void Input::pair_style()
force->pair->settings(narg-1,&arg[1]);
return;
}
- force->create_pair(arg[0]);
+ force->create_pair(arg[0],lmp->asuffix);
if (force->pair) force->pair->settings(narg-1,&arg[1]);
}
diff --git a/src/input.h b/src/input.h
index cbd27881ef..2aac7e9646 100644
--- a/src/input.h
+++ b/src/input.h
@@ -61,7 +61,8 @@ class Input : protected Pointers {
void shell();
void variable_command();
- void angle_coeff(); // LAMMPS commands
+ void accelerator(); // LAMMPS commands
+ void angle_coeff();
void angle_style();
void atom_modify();
void atom_style();
From 0d4486cae801405a8ecce01a84bc67b31e2f6ae1 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:53:47 +0000
Subject: [PATCH 31/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6158
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom_vec_atomic.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/atom_vec_atomic.h b/src/atom_vec_atomic.h
index 187f4c22f8..177038729b 100644
--- a/src/atom_vec_atomic.h
+++ b/src/atom_vec_atomic.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class AtomVecAtomic : public AtomVec {
public:
AtomVecAtomic(class LAMMPS *, int, char **);
- ~AtomVecAtomic() {}
+ virtual ~AtomVecAtomic() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
@@ -50,7 +50,7 @@ class AtomVecAtomic : public AtomVec {
void data_atom(double *, int, char **);
bigint memory_usage();
- private:
+ protected:
int *tag,*type,*mask,*image;
double **x,**v,**f;
};
From 14397987db0b4520cb63067bf9abff573ebecacf Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:53:55 +0000
Subject: [PATCH 32/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6159
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom_vec_charge.h | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/src/atom_vec_charge.h b/src/atom_vec_charge.h
index a19d7d6cde..cc37810598 100644
--- a/src/atom_vec_charge.h
+++ b/src/atom_vec_charge.h
@@ -27,6 +27,7 @@ namespace LAMMPS_NS {
class AtomVecCharge : public AtomVec {
public:
AtomVecCharge(class LAMMPS *, int, char **);
+ virtual ~AtomVecCharge() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
@@ -52,7 +53,7 @@ class AtomVecCharge : public AtomVec {
int data_atom_hybrid(int, char **);
bigint memory_usage();
- private:
+ protected:
int *tag,*type,*mask,*image;
double **x,**v,**f;
double *q;
From 10136a88216287e0049f917c7439cca9606bc546 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:54:13 +0000
Subject: [PATCH 33/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6160
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/MOLECULE/atom_vec_angle.h | 3 ++-
src/MOLECULE/atom_vec_full.h | 3 ++-
2 files changed, 4 insertions(+), 2 deletions(-)
diff --git a/src/MOLECULE/atom_vec_angle.h b/src/MOLECULE/atom_vec_angle.h
index daf18f2198..a9f35573ee 100644
--- a/src/MOLECULE/atom_vec_angle.h
+++ b/src/MOLECULE/atom_vec_angle.h
@@ -27,6 +27,7 @@ namespace LAMMPS_NS {
class AtomVecAngle : public AtomVec {
public:
AtomVecAngle(class LAMMPS *, int, char **);
+ virtual ~AtomVecAngle() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
@@ -52,7 +53,7 @@ class AtomVecAngle : public AtomVec {
int data_atom_hybrid(int, char **);
bigint memory_usage();
- private:
+ protected:
int *tag,*type,*mask,*image;
double **x,**v,**f;
int *molecule;
diff --git a/src/MOLECULE/atom_vec_full.h b/src/MOLECULE/atom_vec_full.h
index 408cb518ff..de68deb86c 100644
--- a/src/MOLECULE/atom_vec_full.h
+++ b/src/MOLECULE/atom_vec_full.h
@@ -27,6 +27,7 @@ namespace LAMMPS_NS {
class AtomVecFull : public AtomVec {
public:
AtomVecFull(class LAMMPS *, int, char **);
+ virtual ~AtomVecFull() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
@@ -52,7 +53,7 @@ class AtomVecFull : public AtomVec {
int data_atom_hybrid(int, char **);
bigint memory_usage();
- private:
+ protected:
int *tag,*type,*mask,*image;
double **x,**v,**f;
double *q;
From cec89647a4c1fd431f4b103ef338fad8009b9d35 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:54:40 +0000
Subject: [PATCH 34/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6161
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/fix_viscous.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/fix_viscous.h b/src/fix_viscous.h
index ccadf4889a..de2c5156e1 100644
--- a/src/fix_viscous.h
+++ b/src/fix_viscous.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class FixViscous : public Fix {
public:
FixViscous(class LAMMPS *, int, char **);
- ~FixViscous();
+ virtual ~FixViscous();
int setmask();
void init();
void setup(int);
@@ -36,7 +36,7 @@ class FixViscous : public Fix {
void post_force_respa(int, int, int);
void min_post_force(int);
- private:
+ protected:
double *gamma;
int nlevels_respa;
};
From 1ca6e02a86c51d57b08ace89db3a61c2f941be0a Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:55:12 +0000
Subject: [PATCH 35/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6162
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/compute_temp_partial.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/compute_temp_partial.h b/src/compute_temp_partial.h
index 83bd6de752..e65c0b2191 100644
--- a/src/compute_temp_partial.h
+++ b/src/compute_temp_partial.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class ComputeTempPartial : public Compute {
public:
ComputeTempPartial(class LAMMPS *, int, char **);
- ~ComputeTempPartial();
+ virtual ~ComputeTempPartial();
void init();
double compute_scalar();
void compute_vector();
@@ -39,7 +39,7 @@ class ComputeTempPartial : public Compute {
void restore_bias_all();
double memory_usage();
- private:
+ protected:
int xflag,yflag,zflag;
int fix_dof;
double tfactor;
From a89bb79df3a6180873c53d6953aa1a1c8714326e Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:55:33 +0000
Subject: [PATCH 36/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6163
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/pair_buck.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/pair_buck.h b/src/pair_buck.h
index 70d571c8b8..c9e13d6cfa 100644
--- a/src/pair_buck.h
+++ b/src/pair_buck.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairBuck : public Pair {
public:
PairBuck(class LAMMPS *);
- ~PairBuck();
+ virtual ~PairBuck();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -39,7 +39,7 @@ class PairBuck : public Pair {
double single(int, int, int, int, double, double, double, double &);
void *extract(char *, int &);
- private:
+ protected:
double cut_global;
double **cut;
double **a,**rho,**c;
From 4b6f064dd72c567a9f429916ef2528f30887d0c5 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:55:58 +0000
Subject: [PATCH 37/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6164
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/compute_pressure.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/compute_pressure.h b/src/compute_pressure.h
index 2e362ab222..cab96ae619 100644
--- a/src/compute_pressure.h
+++ b/src/compute_pressure.h
@@ -27,13 +27,13 @@ namespace LAMMPS_NS {
class ComputePressure : public Compute {
public:
ComputePressure(class LAMMPS *, int, char **);
- ~ComputePressure();
+ virtual ~ComputePressure();
void init();
double compute_scalar();
void compute_vector();
void reset_extra_compute_fix(char *);
- private:
+ protected:
double boltz,nktv2p,inv_volume;
int nvirial,dimension;
double **vptr;
From c41b9fa472e6c3cd7f5418d974754e7f3880e21a Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:56:34 +0000
Subject: [PATCH 38/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6165
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/GRANULAR/fix_pour.h | 1 +
1 file changed, 1 insertion(+)
diff --git a/src/GRANULAR/fix_pour.h b/src/GRANULAR/fix_pour.h
index 5d03a6dc40..d279f54046 100644
--- a/src/GRANULAR/fix_pour.h
+++ b/src/GRANULAR/fix_pour.h
@@ -28,6 +28,7 @@ class FixPour : public Fix {
friend class PairGranHertzHistory;
friend class PairGranHooke;
friend class PairGranHookeHistory;
+ friend class PairGranHookeCuda;
public:
FixPour(class LAMMPS *, int, char **);
From 7bef76d7f2373141ab97c724390723c8fe50ecc6 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:57:20 +0000
Subject: [PATCH 39/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6166
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/KSPACE/pair_born_coul_long.h | 4 ++--
src/KSPACE/pair_buck_coul_long.h | 4 ++--
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/src/KSPACE/pair_born_coul_long.h b/src/KSPACE/pair_born_coul_long.h
index dfd77c4bb6..9cbeb9deb0 100644
--- a/src/KSPACE/pair_born_coul_long.h
+++ b/src/KSPACE/pair_born_coul_long.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairBornCoulLong : public Pair {
public:
PairBornCoulLong(class LAMMPS *);
- ~PairBornCoulLong();
+ virtual ~PairBornCoulLong();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -40,7 +40,7 @@ class PairBornCoulLong : public Pair {
double single(int, int, int, int, double, double, double, double &);
void *extract(char *, int &);
- private:
+ protected:
double cut_lj_global;
double **cut_lj,**cut_ljsq;
double cut_coul,cut_coulsq;
diff --git a/src/KSPACE/pair_buck_coul_long.h b/src/KSPACE/pair_buck_coul_long.h
index ada634a11b..44b82d99de 100644
--- a/src/KSPACE/pair_buck_coul_long.h
+++ b/src/KSPACE/pair_buck_coul_long.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairBuckCoulLong : public Pair {
public:
PairBuckCoulLong(class LAMMPS *);
- ~PairBuckCoulLong();
+ virtual ~PairBuckCoulLong();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -40,7 +40,7 @@ class PairBuckCoulLong : public Pair {
double single(int, int, int, int, double, double, double, double &);
void *extract(char *, int &);
- private:
+ protected:
double cut_lj_global;
double **cut_lj,**cut_ljsq;
double cut_coul,cut_coulsq;
From dbf6c1542ec6ca3cbc444d410bd96ed9e2d4e53b Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:58:03 +0000
Subject: [PATCH 40/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6167
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/CLASS2/pair_lj_class2.h | 4 ++--
src/CLASS2/pair_lj_class2_coul_cut.h | 4 ++--
src/CLASS2/pair_lj_class2_coul_long.h | 4 ++--
3 files changed, 6 insertions(+), 6 deletions(-)
diff --git a/src/CLASS2/pair_lj_class2.h b/src/CLASS2/pair_lj_class2.h
index f2c5a22841..764b8b8c92 100644
--- a/src/CLASS2/pair_lj_class2.h
+++ b/src/CLASS2/pair_lj_class2.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairLJClass2 : public Pair {
public:
PairLJClass2(class LAMMPS *);
- ~PairLJClass2();
+ virtual ~PairLJClass2();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -38,7 +38,7 @@ class PairLJClass2 : public Pair {
void read_restart_settings(FILE *);
double single(int, int, int, int, double, double, double, double &);
- private:
+ protected:
double cut_global;
double **cut;
double **epsilon,**sigma;
diff --git a/src/CLASS2/pair_lj_class2_coul_cut.h b/src/CLASS2/pair_lj_class2_coul_cut.h
index 0fd36104dc..12fabc3f1a 100644
--- a/src/CLASS2/pair_lj_class2_coul_cut.h
+++ b/src/CLASS2/pair_lj_class2_coul_cut.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairLJClass2CoulCut : public Pair {
public:
PairLJClass2CoulCut(class LAMMPS *);
- ~PairLJClass2CoulCut();
+ virtual ~PairLJClass2CoulCut();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -39,7 +39,7 @@ class PairLJClass2CoulCut : public Pair {
void read_restart_settings(FILE *);
double single(int, int, int, int, double, double, double, double &);
- private:
+ protected:
double cut_lj_global,cut_coul_global;
double **cut_lj,**cut_ljsq;
double **cut_coul,**cut_coulsq;
diff --git a/src/CLASS2/pair_lj_class2_coul_long.h b/src/CLASS2/pair_lj_class2_coul_long.h
index 154d67412c..d1b19f84cf 100644
--- a/src/CLASS2/pair_lj_class2_coul_long.h
+++ b/src/CLASS2/pair_lj_class2_coul_long.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairLJClass2CoulLong : public Pair {
public:
PairLJClass2CoulLong(class LAMMPS *);
- ~PairLJClass2CoulLong();
+ virtual ~PairLJClass2CoulLong();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -40,7 +40,7 @@ class PairLJClass2CoulLong : public Pair {
double single(int, int, int, int, double, double, double, double &);
void *extract(char *, int &);
- private:
+ protected:
double cut_lj_global;
double **cut_lj,**cut_ljsq;
double cut_coul,cut_coulsq;
From 478b1f979469463c24ccdcd046af8094c1591d23 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:58:41 +0000
Subject: [PATCH 41/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6168
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/fix_nve.h | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/src/fix_nve.h b/src/fix_nve.h
index f894b8dc6c..8ada964608 100644
--- a/src/fix_nve.h
+++ b/src/fix_nve.h
@@ -31,9 +31,9 @@ class FixNVE : public Fix {
virtual void init();
virtual void initial_integrate(int);
virtual void final_integrate();
- void initial_integrate_respa(int, int, int);
- void final_integrate_respa(int, int);
- void reset_dt();
+ virtual void initial_integrate_respa(int, int, int);
+ virtual void final_integrate_respa(int, int);
+ virtual void reset_dt();
protected:
double dtv,dtf;
From 064bf3b8f7e014b5bfe482092b6bd51360089f9b Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:59:07 +0000
Subject: [PATCH 42/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6169
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/pair_lj_smooth.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/pair_lj_smooth.h b/src/pair_lj_smooth.h
index 1e0faf9c83..6b2fc150bb 100644
--- a/src/pair_lj_smooth.h
+++ b/src/pair_lj_smooth.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairLJSmooth : public Pair {
public:
PairLJSmooth(class LAMMPS *);
- ~PairLJSmooth();
+ virtual ~PairLJSmooth();
void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
@@ -38,7 +38,7 @@ class PairLJSmooth : public Pair {
void read_restart_settings(FILE *);
double single(int, int, int, int, double, double, double, double &);
- private:
+ protected:
double cut_inner_global,cut_global;
double **cut,**cut_inner,**cut_inner_sq;
double **epsilon,**sigma;
From 8664dffaa61bf564cc251eaaa0f597e6ef123913 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 15:59:31 +0000
Subject: [PATCH 43/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6170
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/fix_wall_reflect.h | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/fix_wall_reflect.h b/src/fix_wall_reflect.h
index 0bf98a33c0..6b87d4039b 100644
--- a/src/fix_wall_reflect.h
+++ b/src/fix_wall_reflect.h
@@ -27,12 +27,12 @@ namespace LAMMPS_NS {
class FixWallReflect : public Fix {
public:
FixWallReflect(class LAMMPS *, int, char **);
- ~FixWallReflect();
+ virtual ~FixWallReflect();
int setmask();
void init();
void post_integrate();
- private:
+ protected:
int nwall;
int wallwhich[6],wallstyle[6];
double coord0[6];
From 4d53daece8967b8cb2d0f1be33b0306c7d89e414 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:00:29 +0000
Subject: [PATCH 44/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6171
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/comm.cpp | 2 +-
src/comm.h | 6 +++---
2 files changed, 4 insertions(+), 4 deletions(-)
diff --git a/src/comm.cpp b/src/comm.cpp
index e97d3f3474..c8c06bbf16 100644
--- a/src/comm.cpp
+++ b/src/comm.cpp
@@ -400,7 +400,7 @@ void Comm::setup()
other per-atom attributes may also be sent via pack/unpack routines
------------------------------------------------------------------------- */
-void Comm::forward_comm()
+void Comm::forward_comm(int dummy)
{
int n;
MPI_Request request;
diff --git a/src/comm.h b/src/comm.h
index 11b7e28b8a..d6934775bf 100644
--- a/src/comm.h
+++ b/src/comm.h
@@ -35,12 +35,12 @@ class Comm : protected Pointers {
int ***grid2proc; // which proc owns i,j,k loc in 3d grid
Comm(class LAMMPS *);
- ~Comm();
+ virtual ~Comm();
void init();
void set_procs(); // setup 3d grid of procs
void setup(); // setup 3d communication pattern
- void forward_comm(); // forward communication of atom coords
+ void forward_comm(int dummy = 0); // forward communication of atom coords
void reverse_comm(); // reverse communication of forces
void exchange(); // move atoms to new procs
void borders(); // setup list of atoms to communicate
@@ -55,7 +55,7 @@ class Comm : protected Pointers {
void set(int, char **); // set communication style
bigint memory_usage();
- private:
+ protected:
int style; // single vs multi-type comm
int nswap; // # of swaps to perform
int need[3]; // procs I need atoms from in each dim
From b02b7636fbf705c53af0da1b5aeaea2be5e8cec0 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:00:49 +0000
Subject: [PATCH 45/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6172
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/lammps.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/lammps.cpp b/src/lammps.cpp
index 9ccbcbde70..4ac4307ddf 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -32,7 +32,7 @@
using namespace LAMMPS_NS;
-enum{NONE,OPT,GPU,USERCUDA};
+enum{NOACCEL,OPT,GPU,USERCUDA};
/* ----------------------------------------------------------------------
start up LAMMPS
From add7cba329c4d2388e3413a597d89f52f6b86f14 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:01:08 +0000
Subject: [PATCH 46/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6173
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/pair_lj_cut_coul_debye.h | 3 ++-
1 file changed, 2 insertions(+), 1 deletion(-)
diff --git a/src/pair_lj_cut_coul_debye.h b/src/pair_lj_cut_coul_debye.h
index f112e5df7e..d008884801 100644
--- a/src/pair_lj_cut_coul_debye.h
+++ b/src/pair_lj_cut_coul_debye.h
@@ -27,13 +27,14 @@ namespace LAMMPS_NS {
class PairLJCutCoulDebye : public PairLJCutCoulCut {
public:
PairLJCutCoulDebye(class LAMMPS *);
+ virtual ~PairLJCutCoulDebye() {}
void compute(int, int);
void settings(int, char **);
void write_restart_settings(FILE *);
void read_restart_settings(FILE *);
double single(int, int, int, int, double, double, double, double &);
- private:
+ protected:
double kappa;
};
From 9573dfcea03b51ae12a12709cbcab143a28d0e6b Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:03:21 +0000
Subject: [PATCH 47/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6174
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/fix.cpp | 1 +
src/fix.h | 1 +
2 files changed, 2 insertions(+)
diff --git a/src/fix.cpp b/src/fix.cpp
index e2c077fbc4..2ed2dc7485 100644
--- a/src/fix.cpp
+++ b/src/fix.cpp
@@ -58,6 +58,7 @@ Fix::Fix(LAMMPS *lmp, int narg, char **arg) : Pointers(lmp)
time_depend = 0;
create_attribute = 0;
restart_pbc = 0;
+ cudable_comm = 0;
scalar_flag = vector_flag = array_flag = 0;
peratom_flag = local_flag = 0;
diff --git a/src/fix.h b/src/fix.h
index ca7008a8b2..b1ede8b67c 100644
--- a/src/fix.h
+++ b/src/fix.h
@@ -41,6 +41,7 @@ class Fix : protected Pointers {
// setting when a new atom is created
int restart_pbc; // 1 if fix moves atoms (except integrate)
// so write_restart must remap to PBC
+ int cudable_comm; // 1 if fix has CUDA-enabled communication
int scalar_flag; // 0/1 if compute_scalar() function exists
int vector_flag; // 0/1 if compute_vector() function exists
From 033b04a9699cad1123c08f10af65897b4164df04 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:05:35 +0000
Subject: [PATCH 48/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6175
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/thermo.cpp | 4 ++++
src/thermo.h | 2 ++
2 files changed, 6 insertions(+)
diff --git a/src/thermo.cpp b/src/thermo.cpp
index 3e642a925e..afb32ac9a2 100644
--- a/src/thermo.cpp
+++ b/src/thermo.cpp
@@ -219,12 +219,16 @@ void Thermo::init()
}
// find current ptr for each Compute ID
+ // cudable = 0 if any compute used by Thermo is non-CUDA
+
+ cudable = 1;
int icompute;
for (i = 0; i < ncompute; i++) {
icompute = modify->find_compute(id_compute[i]);
if (icompute < 0) error->all("Could not find thermo compute ID");
computes[i] = modify->compute[icompute];
+ cudable = cudable && computes[i]->cudable;
}
// find current ptr for each Fix ID
diff --git a/src/thermo.h b/src/thermo.h
index 9b9133c6a6..49e2f79e57 100644
--- a/src/thermo.h
+++ b/src/thermo.h
@@ -26,6 +26,7 @@ class Thermo : protected Pointers {
char *style;
int normflag; // 0 if do not normalize by atoms, 1 if normalize
int modified; // 1 if thermo_modify has been used, else 0
+ int cudable; // 1 if all computes used are cudable
Thermo(class LAMMPS *, int, char **);
~Thermo();
@@ -73,6 +74,7 @@ class Thermo : protected Pointers {
int *field2index; // which compute,fix,variable calcs this field
int *argindex1; // indices into compute,fix scalar,vector
int *argindex2;
+
// data for keyword-specific Compute objects
// index = where they are in computes list
// id = ID of Compute objects
From ae3bb4a81ecdb828cce9ddf9783a6453a102b371 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:06:05 +0000
Subject: [PATCH 49/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6176
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/output.cpp | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/src/output.cpp b/src/output.cpp
index 87e85286be..a7339c1f58 100644
--- a/src/output.cpp
+++ b/src/output.cpp
@@ -52,18 +52,18 @@ Output::Output(LAMMPS *lmp) : Pointers(lmp)
newarg[0] = (char *) "thermo_temp";
newarg[1] = (char *) "all";
newarg[2] = (char *) "temp";
- modify->add_compute(3,newarg);
+ modify->add_compute(3,newarg,lmp->asuffix);
newarg[0] = (char *) "thermo_press";
newarg[1] = (char *) "all";
newarg[2] = (char *) "pressure";
newarg[3] = (char *) "thermo_temp";
- modify->add_compute(4,newarg);
+ modify->add_compute(4,newarg,lmp->asuffix);
newarg[0] = (char *) "thermo_pe";
newarg[1] = (char *) "all";
newarg[2] = (char *) "pe";
- modify->add_compute(3,newarg);
+ modify->add_compute(3,newarg,lmp->asuffix);
delete [] newarg;
From fce3d4b87529a9e7e929c4c1dde533f9a92af253 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:08:48 +0000
Subject: [PATCH 50/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6177
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/lammps.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/lammps.cpp b/src/lammps.cpp
index 4ac4307ddf..bc70ebc9e4 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -57,7 +57,7 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
int inflag = 0;
int screenflag = 0;
int logflag = 0;
- accelerator = NONE;
+ accelerator = NOACCEL;
cuda = NULL;
asuffix = NULL;
From 9879ce34aa53c073bad20e45eb8e129738cf1210 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:08:57 +0000
Subject: [PATCH 51/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6178
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/update.cpp | 21 ++++++++++++++++-----
1 file changed, 16 insertions(+), 5 deletions(-)
diff --git a/src/update.cpp b/src/update.cpp
index c6b135eec6..6750343cda 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -25,11 +25,14 @@
#include "region.h"
#include "compute.h"
#include "output.h"
+#include "accelerator.h"
#include "memory.h"
#include "error.h"
using namespace LAMMPS_NS;
+enum{NOACCEL,OPT,GPU,USERCUDA}; // same as lammps.cpp
+
/* ---------------------------------------------------------------------- */
Update::Update(LAMMPS *lmp) : Pointers(lmp)
@@ -53,11 +56,19 @@ Update::Update(LAMMPS *lmp) : Pointers(lmp)
unit_style = NULL;
set_units("lj");
- str = (char *) "verlet";
- n = strlen(str) + 1;
- integrate_style = new char[n];
- strcpy(integrate_style,str);
- integrate = new Verlet(lmp,0,NULL);
+ if (lmp->accelerator == USERCUDA) {
+ str = (char *) "verlet/cuda";
+ n = strlen(str) + 1;
+ integrate_style = new char[n];
+ strcpy(integrate_style,str);
+ integrate = new VerletCuda(lmp,0,NULL);
+ } else {
+ str = (char *) "verlet";
+ n = strlen(str) + 1;
+ integrate_style = new char[n];
+ strcpy(integrate_style,str);
+ integrate = new Verlet(lmp,0,NULL);
+ }
str = (char *) "cg";
n = strlen(str) + 1;
From f99f923320499025b182138d8b7356bc980fba5b Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:16:34 +0000
Subject: [PATCH 52/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6179
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/update.cpp | 32 ++++++++++++++++++++++++++++----
1 file changed, 28 insertions(+), 4 deletions(-)
diff --git a/src/update.cpp b/src/update.cpp
index 6750343cda..fb04ecb3b7 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -202,22 +202,46 @@ void Update::set_units(const char *style)
/* ---------------------------------------------------------------------- */
-void Update::create_integrate(int narg, char **arg)
+void Update::create_integrate(int narg, char **arg, char *suffix)
{
if (narg < 1) error->all("Illegal run_style command");
delete [] integrate_style;
delete integrate;
- if (0) return; // dummy line to enable else-if macro expansion
+ // create the Integrate, first with suffix appended
+
+ int success = 0;
+
+ if (suffix) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",arg[0],suffix);
+ success = 1;
+
+ if (0) return;
#define INTEGRATE_CLASS
#define IntegrateStyle(key,Class) \
- else if (strcmp(arg[0],#key) == 0) integrate = new Class(lmp,narg-1,&arg[1]);
+ else if (strcmp(estyle,#key) == 0) \
+ integrate = new Class(lmp,narg-1,&arg[1]);
#include "style_integrate.h"
#undef INTEGRATE_CLASS
- else error->all("Illegal run_style command");
+ else success = 0;
+ }
+
+ if (!success) {
+ if (0) return;
+
+#define INTEGRATE_CLASS
+#define IntegrateStyle(key,Class) \
+ else if (strcmp(arg[0],#key) == 0) \
+ integrate = new Class(lmp,narg-1,&arg[1]);
+#include "style_integrate.h"
+#undef INTEGRATE_CLASS
+
+ else error->all("Illegal run_style command");
+ }
int n = strlen(arg[0]) + 1;
integrate_style = new char[n];
From e5693fb5db2da1cd323bfe72d506ba239052915a Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:16:48 +0000
Subject: [PATCH 53/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6180
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/input.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/input.cpp b/src/input.cpp
index de4dcb97e9..ce58398f15 100644
--- a/src/input.cpp
+++ b/src/input.cpp
@@ -1207,7 +1207,7 @@ void Input::run_style()
{
if (domain->box_exist == 0)
error->all("Run_style command before simulation box is defined");
- update->create_integrate(narg,arg);
+ update->create_integrate(narg,arg,lmp->asuffix);
}
/* ---------------------------------------------------------------------- */
From 8ce1721b6a1eed44cdcba3bf4b4e516a98f00a0f Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:21:51 +0000
Subject: [PATCH 54/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6181
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom.cpp | 70 ++++++++++++++++++++++++++++++++++++++++++++--------
src/atom.h | 7 ++++--
2 files changed, 65 insertions(+), 12 deletions(-)
diff --git a/src/atom.cpp b/src/atom.cpp
index 806475b3bc..bef3d991e8 100644
--- a/src/atom.cpp
+++ b/src/atom.cpp
@@ -12,6 +12,7 @@
------------------------------------------------------------------------- */
#include "mpi.h"
+#include "math.h"
#include "stdio.h"
#include "stdlib.h"
#include "string.h"
@@ -30,6 +31,7 @@
#include "update.h"
#include "domain.h"
#include "group.h"
+#include "accelerator.h"
#include "memory.h"
#include "error.h"
@@ -38,6 +40,9 @@ using namespace LAMMPS_NS;
#define DELTA 1
#define DELTA_MEMSTR 1024
#define EPSILON 1.0e-6
+#define CUDA_CHUNK 3000
+
+enum{NOACCEL,OPT,GPU,USERCUDA}; // same as lammps.cpp
#define MIN(A,B) ((A) < (B)) ? (A) : (B)
#define MAX(A,B) ((A) > (B)) ? (A) : (B)
@@ -241,7 +246,7 @@ void Atom::settings(Atom *old)
called from input script, restart file, replicate
------------------------------------------------------------------------- */
-void Atom::create_avec(const char *style, int narg, char **arg)
+void Atom::create_avec(const char *style, int narg, char **arg, char *suffix)
{
delete [] atom_style;
if (avec) delete avec;
@@ -256,7 +261,7 @@ void Atom::create_avec(const char *style, int narg, char **arg)
rmass_flag = radius_flag = omega_flag = torque_flag = angmom_flag = 0;
vfrac_flag = spin_flag = eradius_flag = ervel_flag = erforce_flag = 0;
- avec = new_avec(style,narg,arg);
+ avec = new_avec(style,narg,arg,suffix);
int n = strlen(style) + 1;
atom_style = new char[n];
strcpy(atom_style,style);
@@ -268,20 +273,42 @@ void Atom::create_avec(const char *style, int narg, char **arg)
}
/* ----------------------------------------------------------------------
- generate an AtomVec class
+ generate an AtomVec class, first with suffix appended
------------------------------------------------------------------------- */
-AtomVec *Atom::new_avec(const char *style, int narg, char **arg)
+AtomVec *Atom::new_avec(const char *style, int narg, char **arg, char *suffix)
{
- if (0) return NULL;
+ int success = 0;
+
+ if (suffix) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",style,suffix);
+ success = 1;
+
+ if (0) return NULL;
#define ATOM_CLASS
#define AtomStyle(key,Class) \
- else if (strcmp(style,#key) == 0) return new Class(lmp,narg,arg);
+ else if (strcmp(estyle,#key) == 0) return new Class(lmp,narg,arg);
+#include "style_atom.h"
+#undef AtomStyle
+#undef ATOM_CLASS
+
+ else success = 0;
+ }
+
+ if (!success) {
+ if (0) return NULL;
+
+#define ATOM_CLASS
+#define AtomStyle(key,Class) \
+ else if (strcmp(style,#key) == 0) return new Class(lmp,narg,arg);
#include "style_atom.h"
#undef ATOM_CLASS
- else error->all("Invalid atom style");
+ else error->all("Invalid atom style");
+ }
+
return NULL;
}
@@ -1298,6 +1325,11 @@ void Atom::sort()
nextsort = (update->ntimestep/sortfreq)*sortfreq + sortfreq;
+ // download data from GPU if necessary
+
+ if (lmp->accelerator == USERCUDA && !lmp->cuda->oncpu)
+ lmp->cuda->downloadAll();
+
// re-setup sort bins if needed
if (domain->box_change) setup_sort_bins();
@@ -1373,6 +1405,11 @@ void Atom::sort()
current[empty] = permute[empty];
}
+ // upload data back to GPU if necessary
+
+ if (lmp->accelerator == USERCUDA && !lmp->cuda->oncpu)
+ lmp->cuda->uploadAll();
+
// sanity check that current = permute
//int flag = 0;
@@ -1389,12 +1426,25 @@ void Atom::sort()
void Atom::setup_sort_bins()
{
- // binsize = user setting or 1/2 of neighbor cutoff
- // neighbor cutoff can be 0.0
+ // binsize = user setting or default
+ // default = 1/2 of neighbor cutoff for non-CUDA
+ // CUDA_CHUNK atoms/proc for CUDA
+ // check if neighbor cutoff = 0.0
double binsize;
if (userbinsize > 0.0) binsize = userbinsize;
- else binsize = 0.5 * neighbor->cutneighmax;
+ else if (lmp->accelerator == USERCUDA) {
+ if (domain->dimension == 3) {
+ double vol = (domain->boxhi[0]-domain->boxlo[0]) *
+ (domain->boxhi[1]-domain->boxlo[1]) *
+ (domain->boxhi[2]-domain->boxlo[2]);
+ binsize = pow(CUDA_CHUNK/natoms*vol,1.0/3.0);
+ } else {
+ double area = (domain->boxhi[0]-domain->boxlo[0]) *
+ (domain->boxhi[1]-domain->boxlo[1]);
+ binsize = pow(CUDA_CHUNK/natoms*area,1.0/2.0);
+ }
+ } else binsize = 0.5 * neighbor->cutneighmax;
if (binsize == 0.0) error->all("Atom sorting has bin size = 0.0");
double bininv = 1.0/binsize;
diff --git a/src/atom.h b/src/atom.h
index c0e5dafa60..8d971a9faf 100644
--- a/src/atom.h
+++ b/src/atom.h
@@ -114,8 +114,8 @@ class Atom : protected Pointers {
~Atom();
void settings(class Atom *);
- void create_avec(const char *, int, char **);
- class AtomVec *new_avec(const char *, int, char **);
+ void create_avec(const char *, int, char **, char *suffix = NULL);
+ class AtomVec *new_avec(const char *, int, char **, char *suffix = NULL);
void init();
void setup();
@@ -155,6 +155,9 @@ class Atom : protected Pointers {
void *extract(char *);
+ inline int* get_map_array() {return map_array;};
+ inline int get_map_size() {return map_tag_max+1;};
+
bigint memory_usage();
int memcheck(const char *);
From 721570617ed3ee72a8c678cc46dac1451683d34e Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:22:00 +0000
Subject: [PATCH 55/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6182
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/update.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/update.h b/src/update.h
index bde3512457..3d2a4430a1 100644
--- a/src/update.h
+++ b/src/update.h
@@ -48,7 +48,7 @@ class Update : protected Pointers {
~Update();
void init();
void set_units(const char *);
- void create_integrate(int, char **);
+ void create_integrate(int, char **, char *);
void create_minimize(int, char **);
void reset_timestep(int, char **);
bigint memory_usage();
From 3d149b4676023ce9c868b70ba49680aad7d0c993 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:22:41 +0000
Subject: [PATCH 56/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6183
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/neigh_request.cpp | 4 ++++
src/neigh_request.h | 4 ++++
2 files changed, 8 insertions(+)
diff --git a/src/neigh_request.cpp b/src/neigh_request.cpp
index 99f7be82c7..82c9f0b267 100644
--- a/src/neigh_request.cpp
+++ b/src/neigh_request.cpp
@@ -50,6 +50,7 @@ NeighRequest::NeighRequest(LAMMPS *lmp) : Pointers(lmp)
special = 1;
dnum = 0;
ghost = 0;
+ cudable = 0;
// default is no copy or skip
@@ -99,6 +100,7 @@ int NeighRequest::identical(NeighRequest *other)
if (special != other->special) same = 0;
if (dnum != other->dnum) same = 0;
if (ghost != other->ghost) same = 0;
+ if (cudable != other->cudable) same = 0;
if (copy != other->copy) same = 0;
if (same_skip(other) == 0) same = 0;
@@ -126,6 +128,7 @@ int NeighRequest::same_kind(NeighRequest *other)
if (half_from_full != other->half_from_full) same = 0;
if (newton != other->newton) same = 0;
if (ghost != other->ghost) same = 0;
+ if (cudable != other->cudable) same = 0;
return same;
}
@@ -174,4 +177,5 @@ void NeighRequest::copy_request(NeighRequest *other)
newton = other->newton;
dnum = other->dnum;
ghost = other->ghost;
+ cudable = other->cudable;
}
diff --git a/src/neigh_request.h b/src/neigh_request.h
index d653307985..afff5966c0 100644
--- a/src/neigh_request.h
+++ b/src/neigh_request.h
@@ -72,6 +72,10 @@ class NeighRequest : protected Pointers {
int ghost;
+ // 1 if neighbor list build will be done on GPU
+
+ int cudable;
+
// set by neighbor and pair_hybrid after all requests are made
// these settings do not change kind value
From db07c036527205d6037c5845a171c8ab40cf6625 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:23:38 +0000
Subject: [PATCH 57/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6184
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/neigh_list.h | 3 +++
1 file changed, 3 insertions(+)
diff --git a/src/neigh_list.h b/src/neigh_list.h
index dfb13b8d9a..63675503cf 100644
--- a/src/neigh_list.h
+++ b/src/neigh_list.h
@@ -73,6 +73,8 @@ class NeighList : protected Pointers {
int **stencil_multi; // list of bin offsets in each stencil
double **distsq_multi; // sq distances to bins in each stencil
+ class CudaNeighList *cuda_list; // CUDA neighbor list
+
NeighList(class LAMMPS *, int);
~NeighList();
void grow(int); // grow maxlocal
@@ -80,6 +82,7 @@ class NeighList : protected Pointers {
int **add_pages(); // add pages to neigh list
void copy_skip_info(int *, int **); // copy skip info from a neigh request
void print_attributes(); // debug routine
+ int get_maxlocal() {return maxatoms;}
bigint memory_usage();
private:
From ab5508e3c876607607f42dee41c0b43d99534685 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 16:35:26 +0000
Subject: [PATCH 58/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6185
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/update.cpp | 2 ++
1 file changed, 2 insertions(+)
diff --git a/src/update.cpp b/src/update.cpp
index fb04ecb3b7..95540fb1e1 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -225,6 +225,7 @@ void Update::create_integrate(int narg, char **arg, char *suffix)
else if (strcmp(estyle,#key) == 0) \
integrate = new Class(lmp,narg-1,&arg[1]);
#include "style_integrate.h"
+#undef IntegrateStyle
#undef INTEGRATE_CLASS
else success = 0;
@@ -238,6 +239,7 @@ void Update::create_integrate(int narg, char **arg, char *suffix)
else if (strcmp(arg[0],#key) == 0) \
integrate = new Class(lmp,narg-1,&arg[1]);
#include "style_integrate.h"
+#undef IntegrateStyle
#undef INTEGRATE_CLASS
else error->all("Illegal run_style command");
From 717139ef0f0cf9bba3bf3153882b783c3813f7f1 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 17:14:41 +0000
Subject: [PATCH 59/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6186
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom.cpp | 36 ++++++++++++++++++++++--------------
src/atom.h | 2 +-
src/force.cpp | 35 +++++++++++++++++++++--------------
src/force.h | 2 +-
4 files changed, 45 insertions(+), 30 deletions(-)
diff --git a/src/atom.cpp b/src/atom.cpp
index bef3d991e8..39b1841955 100644
--- a/src/atom.cpp
+++ b/src/atom.cpp
@@ -261,10 +261,20 @@ void Atom::create_avec(const char *style, int narg, char **arg, char *suffix)
rmass_flag = radius_flag = omega_flag = torque_flag = angmom_flag = 0;
vfrac_flag = spin_flag = eradius_flag = ervel_flag = erforce_flag = 0;
- avec = new_avec(style,narg,arg,suffix);
- int n = strlen(style) + 1;
- atom_style = new char[n];
- strcpy(atom_style,style);
+ int sflag;
+ avec = new_avec(style,narg,arg,suffix,sflag);
+
+ if (sflag) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",style,suffix);
+ int n = strlen(estyle) + 1;
+ atom_style = new char[n];
+ strcpy(atom_style,estyle);
+ } else {
+ int n = strlen(style) + 1;
+ atom_style = new char[n];
+ strcpy(atom_style,style);
+ }
// if molecular system, default is to have array map
@@ -276,14 +286,13 @@ void Atom::create_avec(const char *style, int narg, char **arg, char *suffix)
generate an AtomVec class, first with suffix appended
------------------------------------------------------------------------- */
-AtomVec *Atom::new_avec(const char *style, int narg, char **arg, char *suffix)
+AtomVec *Atom::new_avec(const char *style, int narg, char **arg,
+ char *suffix, int &sflag)
{
- int success = 0;
-
if (suffix) {
+ sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
- success = 1;
if (0) return NULL;
@@ -294,20 +303,19 @@ AtomVec *Atom::new_avec(const char *style, int narg, char **arg, char *suffix)
#undef AtomStyle
#undef ATOM_CLASS
- else success = 0;
}
- if (!success) {
- if (0) return NULL;
+ sflag = 0;
+
+ if (0) return NULL;
#define ATOM_CLASS
#define AtomStyle(key,Class) \
- else if (strcmp(style,#key) == 0) return new Class(lmp,narg,arg);
+ else if (strcmp(style,#key) == 0) return new Class(lmp,narg,arg);
#include "style_atom.h"
#undef ATOM_CLASS
- else error->all("Invalid atom style");
- }
+ else error->all("Invalid atom style");
return NULL;
}
diff --git a/src/atom.h b/src/atom.h
index 8d971a9faf..e4972b27e6 100644
--- a/src/atom.h
+++ b/src/atom.h
@@ -115,7 +115,7 @@ class Atom : protected Pointers {
void settings(class Atom *);
void create_avec(const char *, int, char **, char *suffix = NULL);
- class AtomVec *new_avec(const char *, int, char **, char *suffix = NULL);
+ class AtomVec *new_avec(const char *, int, char **, char *, int &);
void init();
void setup();
diff --git a/src/force.cpp b/src/force.cpp
index 867101dac4..3c6e280f64 100644
--- a/src/force.cpp
+++ b/src/force.cpp
@@ -122,24 +122,32 @@ void Force::create_pair(const char *style, char *suffix)
delete [] pair_style;
if (pair) delete pair;
- pair = new_pair(style,suffix);
- int n = strlen(style) + 1;
- pair_style = new char[n];
- strcpy(pair_style,style);
+ int sflag;
+ pair = new_pair(style,suffix,sflag);
+
+ if (sflag) {
+ char estyle[256];
+ sprintf(estyle,"%s/%s",style,suffix);
+ int n = strlen(estyle) + 1;
+ pair_style = new char[n];
+ strcpy(pair_style,estyle);
+ } else {
+ int n = strlen(style) + 1;
+ pair_style = new char[n];
+ strcpy(pair_style,style);
+ }
}
/* ----------------------------------------------------------------------
generate a pair class, first with suffix appended
------------------------------------------------------------------------- */
-Pair *Force::new_pair(const char *style, char *suffix)
+Pair *Force::new_pair(const char *style, char *suffix, int &sflag)
{
- int success = 0;
-
if (suffix) {
+ sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
- success = 1;
if (0) return NULL;
@@ -150,20 +158,19 @@ Pair *Force::new_pair(const char *style, char *suffix)
#undef PairStyle
#undef PAIR_CLASS
- else success = 0;
}
- if (!success) {
- if (strcmp(style,"none") == 0) return NULL;
+ sflag = 0;
+
+ if (strcmp(style,"none") == 0) return NULL;
#define PAIR_CLASS
#define PairStyle(key,Class) \
- else if (strcmp(style,#key) == 0) return new Class(lmp);
+ else if (strcmp(style,#key) == 0) return new Class(lmp);
#include "style_pair.h"
#undef PAIR_CLASS
- else error->all("Invalid pair style");
- }
+ else error->all("Invalid pair style");
return NULL;
}
diff --git a/src/force.h b/src/force.h
index d8b4a20136..9f20b4c66c 100644
--- a/src/force.h
+++ b/src/force.h
@@ -65,7 +65,7 @@ class Force : protected Pointers {
void init();
void create_pair(const char *, char *suffix = NULL);
- class Pair *new_pair(const char *, char *suffix = NULL);
+ class Pair *new_pair(const char *, char *, int &);
class Pair *pair_match(const char *, int);
void create_bond(const char *);
From b40cd20066c75e9cbbd71db94b6bb3e03f0b164e Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 17:14:51 +0000
Subject: [PATCH 60/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6187
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom_vec_hybrid.cpp | 4 ++--
src/pair_hybrid.cpp | 8 +++++---
2 files changed, 7 insertions(+), 5 deletions(-)
diff --git a/src/atom_vec_hybrid.cpp b/src/atom_vec_hybrid.cpp
index a7f48066f0..03c0c501ea 100644
--- a/src/atom_vec_hybrid.cpp
+++ b/src/atom_vec_hybrid.cpp
@@ -34,7 +34,7 @@ using namespace LAMMPS_NS;
AtomVecHybrid::AtomVecHybrid(LAMMPS *lmp, int narg, char **arg) :
AtomVec(lmp, narg, arg)
{
- int i,k;
+ int i,k,dummy;
if (narg < 1) error->all("Illegal atom_style command");
@@ -50,7 +50,7 @@ AtomVecHybrid::AtomVecHybrid(LAMMPS *lmp, int narg, char **arg) :
error->all("Atom style hybrid cannot use same atom style twice");
if (strcmp(arg[i],"hybrid") == 0)
error->all("Atom style hybrid cannot have hybrid as an argument");
- styles[i] = atom->new_avec(arg[i],0,NULL);
+ styles[i] = atom->new_avec(arg[i],0,NULL,NULL,dummy);
keywords[i] = new char[strlen(arg[i])+1];
strcpy(keywords[i],arg[i]);
}
diff --git a/src/pair_hybrid.cpp b/src/pair_hybrid.cpp
index a3d0dafd32..5fd45a8f60 100644
--- a/src/pair_hybrid.cpp
+++ b/src/pair_hybrid.cpp
@@ -227,6 +227,8 @@ void PairHybrid::settings(int narg, char **arg)
// exception is 1st arg of reax/c style, which is non-numeric
// need a better way to skip these exceptions
+ int dummy;
+
nstyles = 0;
i = 0;
while (i < narg) {
@@ -237,7 +239,7 @@ void PairHybrid::settings(int narg, char **arg)
error->all("Pair style hybrid cannot have hybrid as an argument");
if (strcmp(arg[i],"none") == 0)
error->all("Pair style hybrid cannot have none as an argument");
- styles[nstyles] = force->new_pair(arg[i]);
+ styles[nstyles] = force->new_pair(arg[i],NULL,dummy);
keywords[nstyles] = new char[strlen(arg[i])+1];
strcpy(keywords[nstyles],arg[i]);
istyle = i;
@@ -574,14 +576,14 @@ void PairHybrid::read_restart(FILE *fp)
// each sub-style is created via new_pair()
// each reads its settings, but no coeff info
- int n;
+ int n,dummy;
for (int m = 0; m < nstyles; m++) {
if (me == 0) fread(&n,sizeof(int),1,fp);
MPI_Bcast(&n,1,MPI_INT,0,world);
keywords[m] = new char[n];
if (me == 0) fread(keywords[m],sizeof(char),n,fp);
MPI_Bcast(keywords[m],n,MPI_CHAR,0,world);
- styles[m] = force->new_pair(keywords[m]);
+ styles[m] = force->new_pair(keywords[m],NULL,dummy);
styles[m]->read_restart_settings(fp);
}
}
From 5a5cfcc81b18e1481957fcef6731f289d96cd1f2 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 17:33:20 +0000
Subject: [PATCH 61/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6188
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/update.cpp | 71 +++++++++++++++++++++++++-------------------------
src/update.h | 4 +++
2 files changed, 40 insertions(+), 35 deletions(-)
diff --git a/src/update.cpp b/src/update.cpp
index 95540fb1e1..cdbdaa5fd7 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -56,25 +56,16 @@ Update::Update(LAMMPS *lmp) : Pointers(lmp)
unit_style = NULL;
set_units("lj");
- if (lmp->accelerator == USERCUDA) {
- str = (char *) "verlet/cuda";
- n = strlen(str) + 1;
- integrate_style = new char[n];
- strcpy(integrate_style,str);
- integrate = new VerletCuda(lmp,0,NULL);
- } else {
- str = (char *) "verlet";
- n = strlen(str) + 1;
- integrate_style = new char[n];
- strcpy(integrate_style,str);
- integrate = new Verlet(lmp,0,NULL);
- }
+ integrate_style = NULL;
+ integrate = NULL;
+ minimize_style = NULL;
+ minimize = NULL;
+
+ str = (char *) "verlet";
+ create_integrate(1,&str,lmp->asuffix);
str = (char *) "cg";
- n = strlen(str) + 1;
- minimize_style = new char[n];
- strcpy(minimize_style,str);
- minimize = new MinCG(lmp);
+ create_minimize(1,&str);
}
/* ---------------------------------------------------------------------- */
@@ -209,45 +200,55 @@ void Update::create_integrate(int narg, char **arg, char *suffix)
delete [] integrate_style;
delete integrate;
- // create the Integrate, first with suffix appended
+ int sflag;
+ new_integrate(arg[0],narg-1,&arg[1],suffix,sflag);
- int success = 0;
-
- if (suffix) {
+ if (sflag) {
char estyle[256];
sprintf(estyle,"%s/%s",arg[0],suffix);
- success = 1;
+ int n = strlen(estyle) + 1;
+ integrate_style = new char[n];
+ strcpy(integrate_style,estyle);
+ } else {
+ int n = strlen(arg[0]) + 1;
+ integrate_style = new char[n];
+ strcpy(integrate_style,arg[0]);
+ }
+}
+
+/* ---------------------------------------------------------------------- */
+
+void Update::new_integrate(char *style, int narg, char **arg,
+ char *suffix, int &sflag)
+{
+ if (suffix) {
+ sflag = 1;
+ char estyle[256];
+ sprintf(estyle,"%s/%s",style,suffix);
if (0) return;
#define INTEGRATE_CLASS
#define IntegrateStyle(key,Class) \
- else if (strcmp(estyle,#key) == 0) \
- integrate = new Class(lmp,narg-1,&arg[1]);
+ else if (strcmp(estyle,#key) == 0) integrate = new Class(lmp,narg,arg);
#include "style_integrate.h"
#undef IntegrateStyle
#undef INTEGRATE_CLASS
- else success = 0;
}
- if (!success) {
- if (0) return;
+ sflag = 0;
+
+ if (0) return;
#define INTEGRATE_CLASS
#define IntegrateStyle(key,Class) \
- else if (strcmp(arg[0],#key) == 0) \
- integrate = new Class(lmp,narg-1,&arg[1]);
+ else if (strcmp(style,#key) == 0) integrate = new Class(lmp,narg,arg);
#include "style_integrate.h"
#undef IntegrateStyle
#undef INTEGRATE_CLASS
- else error->all("Illegal run_style command");
- }
-
- int n = strlen(arg[0]) + 1;
- integrate_style = new char[n];
- strcpy(integrate_style,arg[0]);
+ else error->all("Illegal integrate style");
}
/* ---------------------------------------------------------------------- */
diff --git a/src/update.h b/src/update.h
index 3d2a4430a1..692f21548c 100644
--- a/src/update.h
+++ b/src/update.h
@@ -52,6 +52,10 @@ class Update : protected Pointers {
void create_minimize(int, char **);
void reset_timestep(int, char **);
bigint memory_usage();
+
+ private:
+ void new_integrate(char *, int, char **, char *, int &);
+
};
}
From 4d79e92cb164537483439605d2c2b68f1a2d0fad Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 20:17:24 +0000
Subject: [PATCH 62/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6189
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/USER-EFF/atom_vec_electron.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/USER-EFF/atom_vec_electron.cpp b/src/USER-EFF/atom_vec_electron.cpp
index e5c47745ab..c0c97d161b 100644
--- a/src/USER-EFF/atom_vec_electron.cpp
+++ b/src/USER-EFF/atom_vec_electron.cpp
@@ -237,7 +237,7 @@ int AtomVecElectron::pack_comm_hybrid(int n, int *list, double *buf)
m = 0;
for (i = 0; i < n; i++) {
j = list[i];
- buf[m++] = eradius[i];
+ buf[m++] = eradius[j];
}
return m;
}
From dffb8b2a8917d5295b4f8d95668e1e6b87c9fe62 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 20:19:11 +0000
Subject: [PATCH 63/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6190
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/atom_vec_sphere.cpp | 4 ++--
1 file changed, 2 insertions(+), 2 deletions(-)
diff --git a/src/atom_vec_sphere.cpp b/src/atom_vec_sphere.cpp
index ab53b79c5d..eee68209dd 100644
--- a/src/atom_vec_sphere.cpp
+++ b/src/atom_vec_sphere.cpp
@@ -374,8 +374,8 @@ int AtomVecSphere::pack_comm_hybrid(int n, int *list, double *buf)
m = 0;
for (i = 0; i < n; i++) {
j = list[i];
- buf[m++] = radius[i];
- buf[m++] = rmass[i];
+ buf[m++] = radius[j];
+ buf[m++] = rmass[j];
}
return m;
}
From 2739d7ee1d72b23fa7a7d9d2fdfab9ce5f635ddc Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 21:31:49 +0000
Subject: [PATCH 64/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6191
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/lammps.cpp | 4 +++-
1 file changed, 3 insertions(+), 1 deletion(-)
diff --git a/src/lammps.cpp b/src/lammps.cpp
index bc70ebc9e4..95067b0a1d 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -334,7 +334,9 @@ LAMMPS::~LAMMPS()
void LAMMPS::create()
{
atom = new Atom(this);
- neighbor = new Neighbor(this);
+
+ if (accelerator == USERCUDA) neighbor = new NeighborCuda(this);
+ else neighbor = new Neighbor(this);
if (accelerator == USERCUDA) comm = new CommCuda(this);
else comm = new Comm(this);
From 06ffbc8f0a0f6f0c118737e47336cb23ec55b6f2 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 21:33:03 +0000
Subject: [PATCH 65/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6192
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/neighbor.h | 6 +++---
1 file changed, 3 insertions(+), 3 deletions(-)
diff --git a/src/neighbor.h b/src/neighbor.h
index 97e0f354c4..607a145a73 100644
--- a/src/neighbor.h
+++ b/src/neighbor.h
@@ -60,7 +60,7 @@ class Neighbor : protected Pointers {
int **improperlist;
Neighbor(class LAMMPS *);
- ~Neighbor();
+ virtual ~Neighbor();
void init();
int request(void *); // another class requests a neighbor list
void print_lists_of_lists(); // debug print out
@@ -73,7 +73,7 @@ class Neighbor : protected Pointers {
void modify_params(int, char**); // modify parameters that control builds
bigint memory_usage();
- private:
+ protected:
int me,nprocs;
int maxatom; // size of atom-based NeighList arrays
@@ -162,7 +162,7 @@ class Neighbor : protected Pointers {
int coord2bin(double *, int &, int &, int&); // ditto
int exclusion(int, int, int, int, int *, int *); // test for pair exclusion
- void choose_build(int, class NeighRequest *);
+ virtual void choose_build(int, class NeighRequest *);
void choose_stencil(int, class NeighRequest *);
// pairwise build functions
From a764a2e141908985eebcef6bbe74a071842576f3 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 21:33:28 +0000
Subject: [PATCH 66/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6193
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/PERI/atom_vec_peri.cpp | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/PERI/atom_vec_peri.cpp b/src/PERI/atom_vec_peri.cpp
index d85853d114..ee36eb6db8 100644
--- a/src/PERI/atom_vec_peri.cpp
+++ b/src/PERI/atom_vec_peri.cpp
@@ -240,7 +240,7 @@ int AtomVecPeri::pack_comm_hybrid(int n, int *list, double *buf)
m = 0;
for (i = 0; i < n; i++) {
j = list[i];
- buf[m++] = s0[i];
+ buf[m++] = s0[j];
}
return m;
}
From f7eefa416ec1fe5f0777f7c1ab7302af0ad27eab Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 22:29:46 +0000
Subject: [PATCH 67/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6194
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/neighbor.cpp | 12 +++++++++---
src/neighbor.h | 21 ++++++++++++---------
2 files changed, 21 insertions(+), 12 deletions(-)
diff --git a/src/neighbor.cpp b/src/neighbor.cpp
index fc4f1125e5..e8ecccc9a0 100644
--- a/src/neighbor.cpp
+++ b/src/neighbor.cpp
@@ -480,11 +480,11 @@ void Neighbor::init()
// fix/compute requests:
// kind of request = half or full, occasional or not doesn't matter
// if request = half and non-skip pair half/respaouter exists,
- // become copy of that list
+ // become copy of that list if cudable flag matches
// if request = full and non-skip pair full exists,
- // become copy of that list
+ // become copy of that list if cudable flag matches
// if request = half and non-skip pair full exists,
- // become half_from_full of that list
+ // become half_from_full of that list if cudable flag matches
// if no matches, do nothing, fix/compute list will be built directly
// ok if parent is copy list
@@ -534,6 +534,8 @@ void Neighbor::init()
if (requests[i]->half && requests[j]->pair &&
requests[j]->skip == 0 && requests[j]->respaouter) break;
}
+ if (j < nlist && requests[j]->cudable != requests[i]->cudable)
+ j = nlist;
if (j < nlist) {
requests[i]->copy = 1;
lists[i]->listcopy = lists[j];
@@ -542,6 +544,8 @@ void Neighbor::init()
if (requests[i]->half && requests[j]->pair &&
requests[j]->skip == 0 && requests[j]->full) break;
}
+ if (j < nlist && requests[j]->cudable != requests[i]->cudable)
+ j = nlist;
if (j < nlist) {
requests[i]->half = 0;
requests[i]->half_from_full = 1;
@@ -553,11 +557,13 @@ void Neighbor::init()
// set ptrs to pair_build and stencil_create functions for each list
// ptrs set to NULL if not set explicitly
+ // also set cudable to 0 if any neigh list request is not cudable
for (i = 0; i < nlist; i++) {
choose_build(i,requests[i]);
if (style != NSQ) choose_stencil(i,requests[i]);
else stencil_create[i] = NULL;
+ if (!requests[i]->cudable) cudable = 0;
}
// set each list's build/grow/stencil/ghost flags based on neigh request
diff --git a/src/neighbor.h b/src/neighbor.h
index 607a145a73..c2fa635d63 100644
--- a/src/neighbor.h
+++ b/src/neighbor.h
@@ -19,6 +19,8 @@
namespace LAMMPS_NS {
class Neighbor : protected Pointers {
+ friend class Cuda;
+
public:
int style; // 0,1,2 = nsq, bin, multi
int every; // build every this many steps
@@ -29,6 +31,7 @@ class Neighbor : protected Pointers {
int oneatom; // max # of neighbors for one atom
int includegroup; // only build pairwise lists for this group
int build_once; // 1 if only build lists once per run
+ int cudable; // GPU <-> CPU communication flag for CUDA
double skin; // skin distance
double cutneighmin; // min neighbor cutoff for all type pairs
@@ -61,15 +64,15 @@ class Neighbor : protected Pointers {
Neighbor(class LAMMPS *);
virtual ~Neighbor();
- void init();
- int request(void *); // another class requests a neighbor list
- void print_lists_of_lists(); // debug print out
- int decide(); // decide whether to build or not
- int check_distance(); // check max distance moved since last build
- void setup_bins(); // setup bins based on box and cutoff
- void build(); // create all neighbor lists (pair,bond)
- void build_one(int); // create a single neighbor list
- void set(int, char **); // set neighbor style and skin distance
+ virtual void init();
+ int request(void *); // another class requests a neighbor list
+ void print_lists_of_lists(); // debug print out
+ int decide(); // decide whether to build or not
+ virtual int check_distance(); // check max distance moved since last build
+ void setup_bins(); // setup bins based on box and cutoff
+ virtual void build(); // create all neighbor lists (pair,bond)
+ void build_one(int); // create a single neighbor list
+ void set(int, char **); // set neighbor style and skin distance
void modify_params(int, char**); // modify parameters that control builds
bigint memory_usage();
From 91b6d45b539b084da5c3eb2ecb81d7c50c8565cb Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 22:46:32 +0000
Subject: [PATCH 68/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6195
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/Section_commands.html | 44 +++++++++++++-------------
doc/Section_commands.txt | 2 ++
doc/Section_start.html | 22 ++++++++++++-
doc/Section_start.txt | 20 ++++++++++++
doc/fix_gpu.html | 66 +++++++++++++++++++++------------------
doc/fix_gpu.txt | 66 +++++++++++++++++++++------------------
6 files changed, 136 insertions(+), 84 deletions(-)
diff --git a/doc/Section_commands.html b/doc/Section_commands.html
index 5f996268de..039367f5ce 100644
--- a/doc/Section_commands.html
+++ b/doc/Section_commands.html
@@ -311,19 +311,20 @@ default LAMMPS build. These dependencies are listed as Restrictions
in the command's documentation.
@@ -336,14 +337,15 @@ of each style or click on the style itself for a full description:
These are fix styles contributed by users, which can be used if
diff --git a/doc/Section_commands.txt b/doc/Section_commands.txt
index 1c58401303..a8bcd7a79b 100644
--- a/doc/Section_commands.txt
+++ b/doc/Section_commands.txt
@@ -307,6 +307,7 @@ included when LAMMPS was built. Not all packages are included in a
default LAMMPS build. These dependencies are listed as Restrictions
in the command's documentation.
+"accelerator"_accelerator.html,
"angle_coeff"_angle_coeff.html,
"angle_style"_angle_style.html,
"atom_modify"_atom_modify.html,
@@ -414,6 +415,7 @@ of each style or click on the style itself for a full description:
"evaporate"_fix_evaporate.html,
"external"_fix_external.html,
"freeze"_fix_freeze.html,
+"gpu"_fix_gpu.html,
"gravity"_fix_gravity.html,
"heat"_fix_heat.html,
"indent"_fix_indent.html,
diff --git a/doc/Section_start.html b/doc/Section_start.html
index e1c61510d2..999bb8b732 100644
--- a/doc/Section_start.html
+++ b/doc/Section_start.html
@@ -787,7 +787,8 @@ more processors or setup a smaller problem.
which may be used in any order. Either the full word or the
one-letter abbreviation can be used:
-- -echo or -e
+
- -accelerator or -a
+
- -echo or -e
- -partition or -p
- -in or -i
- -log or -l
@@ -800,6 +801,25 @@ one-letter abbreviation can be used:
Here are the details on the options:
+-accelerator style
+
+Use accelerated variants of various styles if they exist. The style
+can be opt or gpu or cuda. The variant styles are part of
+optional packages that LAMMPS can be built with, as described above in
+Section 2.3. The "opt" style corrsponds to the OPT package,
+the "gpu" style to the GPU package, and the "cuda" style to the
+USER-CUDA package. For example, all of the packages provide a
+pair_style lj/cut variant, with styles lj/cut/opt or
+lj/cut/gpu or lj/cut/cuda. These styles can be specified explicitly
+in your input script, e.g. pair_style lj/cut/gpu. But if the
+-accelerator switch is used, you do not need to modify your input
+script. The accelerator suffix (opt,gpu,cuda) is automatically
+appended when the style is created. The suffix is applied to atom,
+pair, fix, compute, and integrate styles. If an accelerated version
+does not exist, the standard version is created. See the
+accelerator command for info on how to turn off
+this option.
+
-echo style
Set the style of command echoing. The style can be none or screen
diff --git a/doc/Section_start.txt b/doc/Section_start.txt
index fdb03c0b1c..57bb2a8975 100644
--- a/doc/Section_start.txt
+++ b/doc/Section_start.txt
@@ -777,6 +777,7 @@ At run time, LAMMPS recognizes several optional command-line switches
which may be used in any order. Either the full word or the
one-letter abbreviation can be used:
+-accelerator or -a
-echo or -e
-partition or -p
-in or -i
@@ -790,6 +791,25 @@ mpirun -np 16 lmp_ibm -var f tmp.out -log my.log -screen none < in.alloy :pre
Here are the details on the options:
+-accelerator style :pre
+
+Use accelerated variants of various styles if they exist. The style
+can be {opt} or {gpu} or {cuda}. The variant styles are part of
+optional packages that LAMMPS can be built with, as described above in
+"Section 2.3"_#2_3. The "opt" style corrsponds to the OPT package,
+the "gpu" style to the GPU package, and the "cuda" style to the
+USER-CUDA package. For example, all of the packages provide a
+"pair_style lj/cut"_pair_lj.html variant, with styles lj/cut/opt or
+lj/cut/gpu or lj/cut/cuda. These styles can be specified explicitly
+in your input script, e.g. pair_style lj/cut/gpu. But if the
+-accelerator switch is used, you do not need to modify your input
+script. The accelerator suffix (opt,gpu,cuda) is automatically
+appended when the style is created. The suffix is applied to atom,
+pair, fix, compute, and integrate styles. If an accelerated version
+does not exist, the standard version is created. See the
+"accelerator"_accelerator.html command for info on how to turn off
+this option.
+
-echo style :pre
Set the style of command echoing. The style can be {none} or {screen}
diff --git a/doc/fix_gpu.html b/doc/fix_gpu.html
index f71a8e8a4a..3560baf21f 100644
--- a/doc/fix_gpu.html
+++ b/doc/fix_gpu.html
@@ -45,39 +45,38 @@ specified for a run or an error will be generated. The fix will not have an
effect on any LAMMPS computations that do not use GPU acceleration, so there
should not be any problems with specifying this fix first in input scripts.
-mode specifies where neighbor list calculations will be performed.
-If mode is force, neighbor list calculation is performed on the
-CPU. If mode is force/neigh, neighbor list calculation is
-performed on the GPU. GPU neighbor list calculation currently cannot be
-used with a triclinic box. GPU neighbor list calculation currently
-cannot be used with hybrid pair styles.
-GPU neighbor lists are not compatible with styles that are not GPU-enabled.
-When a non-GPU enabled style requires a neighbor list, it will also be
-built using CPU routines. In these cases, it will typically be more efficient
-to only use CPU neighbor list builds.
+
The mode setting specifies where neighbor list calculations will be
+performed. If mode is force, neighbor list calculation is performed
+on the CPU. If mode is force/neigh, neighbor list calculation is
+performed on the GPU. GPU neighbor list calculation currently cannot
+be used with a triclinic box. GPU neighbor list calculation currently
+cannot be used with hybrid pair styles. GPU
+neighbor lists are not compatible with styles that are not
+GPU-enabled. When a non-GPU enabled style requires a neighbor list,
+it will also be built using CPU routines. In these cases, it will
+typically be more efficient to only use CPU neighbor list builds.
-first and last specify the GPUs that will be used for simulation.
-On each node, the GPU IDs in the inclusive range from first to last will
-be used.
+
The first and last settings specify the GPUs that will be used for
+simulation. On each node, the GPU IDs in the inclusive range from
+first to last will be used.
-split can be used for load balancing force calculation work between
-CPU and GPU cores in GPU-enabled pair styles. If 0<split<1.0,
-a fixed fraction of particles is offloaded to the GPU while force calculation
-for the other particles occurs simulataneously on the CPU. If split<0,
-the optimal fraction (based on CPU and GPU timings) is calculated
-every 25 timesteps. If split=1.0, all force calculations for
-GPU accelerated pair styles are performed
-on the GPU. In this case, hybrid,
-bond, angle,
-dihedral, improper,
-and long-range calculations can be performed on the CPU
-while the GPU is performing force calculations for the GPU-enabled pair
-style.
+
The split setting can be used for load balancing force calculation
+work between CPU and GPU cores in GPU-enabled pair styles. If
+0<split<1.0, a fixed fraction of particles is offloaded to the GPU
+while force calculation for the other particles occurs simulataneously
+on the CPU. If split<0, the optimal fraction (based on CPU and GPU
+timings) is calculated every 25 timesteps. If split=1.0, all force
+calculations for GPU accelerated pair styles are performed on the
+GPU. In this case, hybrid, bond,
+angle, dihedral,
+improper, and long-range
+calculations can be performed on the CPU while the GPU is performing
+force calculations for the GPU-enabled pair style.
-In order to use GPU acceleration, a GPU enabled style must be
-selected in the input script in addition to this fix. Currently,
-this is limited to a few pair styles and
-the PPPM kspace style.
+
In order to use GPU acceleration, a GPU enabled style must be selected
+in the input script in addition to this fix. Currently, this is
+limited to a few pair styles and the PPPM kspace
+style.
More details about these settings and various possible hardware
configuration are in this section of the
@@ -85,6 +84,10 @@ manual.
Restart, fix_modify, output, run start/stop, minimize info:
+This fix is part of the "gpu" package. It is only enabled if LAMMPS
+was built with that package. See the Making
+LAMMPS section for more info.
+
No information about this fix is written to binary restart
files. None of the fix_modify options
are relevant to this fix.
@@ -98,7 +101,8 @@ the run command.
mode should not be used with a triclinic box or hybrid
pair styles.
-split must be positive when using hybrid pair styles.
+
The split setting must be positive when using
+hybrid pair styles.
Currently, group-ID must be all.
diff --git a/doc/fix_gpu.txt b/doc/fix_gpu.txt
index df8fbadb8f..e4cd41f1de 100644
--- a/doc/fix_gpu.txt
+++ b/doc/fix_gpu.txt
@@ -36,39 +36,38 @@ specified for a run or an error will be generated. The fix will not have an
effect on any LAMMPS computations that do not use GPU acceleration, so there
should not be any problems with specifying this fix first in input scripts.
-{mode} specifies where neighbor list calculations will be performed.
-If {mode} is force, neighbor list calculation is performed on the
-CPU. If {mode} is force/neigh, neighbor list calculation is
-performed on the GPU. GPU neighbor list calculation currently cannot be
-used with a triclinic box. GPU neighbor list calculation currently
-cannot be used with "hybrid"_pair_hybrid.html pair styles.
-GPU neighbor lists are not compatible with styles that are not GPU-enabled.
-When a non-GPU enabled style requires a neighbor list, it will also be
-built using CPU routines. In these cases, it will typically be more efficient
-to only use CPU neighbor list builds.
+The {mode} setting specifies where neighbor list calculations will be
+performed. If {mode} is force, neighbor list calculation is performed
+on the CPU. If {mode} is force/neigh, neighbor list calculation is
+performed on the GPU. GPU neighbor list calculation currently cannot
+be used with a triclinic box. GPU neighbor list calculation currently
+cannot be used with "hybrid"_pair_hybrid.html pair styles. GPU
+neighbor lists are not compatible with styles that are not
+GPU-enabled. When a non-GPU enabled style requires a neighbor list,
+it will also be built using CPU routines. In these cases, it will
+typically be more efficient to only use CPU neighbor list builds.
-{first} and {last} specify the GPUs that will be used for simulation.
-On each node, the GPU IDs in the inclusive range from {first} to {last} will
-be used.
+The {first} and {last} settings specify the GPUs that will be used for
+simulation. On each node, the GPU IDs in the inclusive range from
+{first} to {last} will be used.
-{split} can be used for load balancing force calculation work between
-CPU and GPU cores in GPU-enabled pair styles. If 0<{split}<1.0,
-a fixed fraction of particles is offloaded to the GPU while force calculation
-for the other particles occurs simulataneously on the CPU. If {split}<0,
-the optimal fraction (based on CPU and GPU timings) is calculated
-every 25 timesteps. If {split}=1.0, all force calculations for
-GPU accelerated pair styles are performed
-on the GPU. In this case, "hybrid"_pair_hybrid.html,
-"bond"_bond_style.html, "angle"_angle_style.html,
-"dihedral"_dihedral_style.html, "improper"_improper_style.html,
-and "long-range"_kspace_style.html calculations can be performed on the CPU
-while the GPU is performing force calculations for the GPU-enabled pair
-style.
+The {split} setting can be used for load balancing force calculation
+work between CPU and GPU cores in GPU-enabled pair styles. If
+0<{split}<1.0, a fixed fraction of particles is offloaded to the GPU
+while force calculation for the other particles occurs simulataneously
+on the CPU. If {split}<0, the optimal fraction (based on CPU and GPU
+timings) is calculated every 25 timesteps. If {split}=1.0, all force
+calculations for GPU accelerated pair styles are performed on the
+GPU. In this case, "hybrid"_pair_hybrid.html, "bond"_bond_style.html,
+"angle"_angle_style.html, "dihedral"_dihedral_style.html,
+"improper"_improper_style.html, and "long-range"_kspace_style.html
+calculations can be performed on the CPU while the GPU is performing
+force calculations for the GPU-enabled pair style.
-In order to use GPU acceleration, a GPU enabled style must be
-selected in the input script in addition to this fix. Currently,
-this is limited to a few "pair styles"_pair_style.html and
-the PPPM "kspace style"_kspace_style.html.
+In order to use GPU acceleration, a GPU enabled style must be selected
+in the input script in addition to this fix. Currently, this is
+limited to a few "pair styles"_pair_style.html and the PPPM "kspace
+style"_kspace_style.html.
More details about these settings and various possible hardware
configuration are in "this section"_Section_start.html#2_8 of the
@@ -76,6 +75,10 @@ manual.
[Restart, fix_modify, output, run start/stop, minimize info:]
+This fix is part of the "gpu" package. It is only enabled if LAMMPS
+was built with that package. See the "Making
+LAMMPS"_Section_start.html#2_3 section for more info.
+
No information about this fix is written to "binary restart
files"_restart.html. None of the "fix_modify"_fix_modify.html options
are relevant to this fix.
@@ -89,7 +92,8 @@ The fix must be the first fix specified for a given run. The force/neigh
{mode} should not be used with a triclinic box or "hybrid"_pair_hybrid.html
pair styles.
-{split} must be positive when using "hybrid"_pair_hybrid.html pair styles.
+The {split} setting must be positive when using
+"hybrid"_pair_hybrid.html pair styles.
Currently, group-ID must be all.
From 3ae193256b4a3e92f3c5a153c8ccfc67a97d41c8 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 22:59:15 +0000
Subject: [PATCH 69/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6196
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/Makefile | 2 +-
src/atom.cpp | 2 +-
src/force.cpp | 2 +-
src/input.cpp | 7 +++++++
src/lammps.cpp | 3 ++-
src/lammps.h | 1 +
src/modify.cpp | 4 ++--
src/update.cpp | 2 +-
8 files changed, 16 insertions(+), 7 deletions(-)
diff --git a/src/Makefile b/src/Makefile
index 621a446289..7473b0e112 100755
--- a/src/Makefile
+++ b/src/Makefile
@@ -17,7 +17,7 @@ PACKAGE = asphere class2 colloid dipole dsmc gpu granular \
kspace manybody meam molecule opt peri poems reax replica \
shock srd xtc
-PACKUSER = user-ackland user-atc user-cd-eam user-cg-cmm user-cuda user-eff \
+PACKUSER = user-ackland user-atc user-cd-eam user-cg-cmm user-eff \
user-ewaldn user-imd user-reaxc user-smd
PACKALL = $(PACKAGE) $(PACKUSER)
diff --git a/src/atom.cpp b/src/atom.cpp
index 39b1841955..9abc972f3c 100644
--- a/src/atom.cpp
+++ b/src/atom.cpp
@@ -289,7 +289,7 @@ void Atom::create_avec(const char *style, int narg, char **arg, char *suffix)
AtomVec *Atom::new_avec(const char *style, int narg, char **arg,
char *suffix, int &sflag)
{
- if (suffix) {
+ if (suffix && lmp->offaccel == 0) {
sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
diff --git a/src/force.cpp b/src/force.cpp
index 3c6e280f64..3364455924 100644
--- a/src/force.cpp
+++ b/src/force.cpp
@@ -144,7 +144,7 @@ void Force::create_pair(const char *style, char *suffix)
Pair *Force::new_pair(const char *style, char *suffix, int &sflag)
{
- if (suffix) {
+ if (suffix && lmp->offaccel == 0) {
sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
diff --git a/src/input.cpp b/src/input.cpp
index ce58398f15..6d38318089 100644
--- a/src/input.cpp
+++ b/src/input.cpp
@@ -811,6 +811,13 @@ void Input::accelerator()
if (strcmp(lmp->asuffix,arg[0]) != 0)
error->all("Accelerator command requires matching command-line -a switch");
+ if (strcmp(arg[0],"off") == 0) {
+ if (narg != 1) error->all("Illegal accelerator command");
+ lmp->offaccel = 1;
+ return;
+ }
+
+ lmp->offaccel = 0;
if (strcmp(arg[0],"cuda") == 0) lmp->cuda->accelerator(narg-1,&arg[1]);
else error->all("Illegal accelerator command");
}
diff --git a/src/lammps.cpp b/src/lammps.cpp
index 95067b0a1d..daf9c5ac50 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -58,8 +58,9 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
int screenflag = 0;
int logflag = 0;
accelerator = NOACCEL;
- cuda = NULL;
asuffix = NULL;
+ offaccel = 0;
+ cuda = NULL;
int iarg = 1;
diff --git a/src/lammps.h b/src/lammps.h
index 095fa1054d..deaa10ebd9 100644
--- a/src/lammps.h
+++ b/src/lammps.h
@@ -44,6 +44,7 @@ class LAMMPS {
int accelerator; // accelerator flag
char *asuffix; // accelerator suffix
+ int offaccel; // 1 if accelerator flag currently disabled
class Cuda *cuda; // CUDA accelerator class
LAMMPS(int, char **, MPI_Comm);
diff --git a/src/modify.cpp b/src/modify.cpp
index 514341fc17..0576cd2c83 100644
--- a/src/modify.cpp
+++ b/src/modify.cpp
@@ -640,7 +640,7 @@ void Modify::add_fix(int narg, char **arg, char *suffix)
int success = 0;
- if (suffix) {
+ if (suffix && lmp->offaccel == 0) {
char estyle[256];
sprintf(estyle,"%s/%s",arg[2],suffix);
success = 1;
@@ -784,7 +784,7 @@ void Modify::add_compute(int narg, char **arg, char *suffix)
int success = 0;
- if (suffix) {
+ if (suffix && lmp->offaccel == 0) {
char estyle[256];
sprintf(estyle,"%s/%s",arg[2],suffix);
success = 1;
diff --git a/src/update.cpp b/src/update.cpp
index cdbdaa5fd7..1c70ad7047 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -221,7 +221,7 @@ void Update::create_integrate(int narg, char **arg, char *suffix)
void Update::new_integrate(char *style, int narg, char **arg,
char *suffix, int &sflag)
{
- if (suffix) {
+ if (suffix && lmp->offaccel == 0) {
sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
From 157ffd67ae38b311478f528a3eabc8c803f4c2c8 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 23:21:12 +0000
Subject: [PATCH 70/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6197
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/accelerator.html | 88 ++++++++++++++++++++++++++++++++++++++++++++
doc/accelerator.txt | 80 ++++++++++++++++++++++++++++++++++++++++
2 files changed, 168 insertions(+)
create mode 100644 doc/accelerator.html
create mode 100644 doc/accelerator.txt
diff --git a/doc/accelerator.html b/doc/accelerator.html
new file mode 100644
index 0000000000..d53813863f
--- /dev/null
+++ b/doc/accelerator.html
@@ -0,0 +1,88 @@
+
+LAMMPS WWW Site - LAMMPS Documentation - LAMMPS Commands
+
+
+
+
+
+
+
+
+
+accelerator command
+
+Syntax:
+
+accelerator style args
+
+
+Examples:
+
+accelerator off
+accelerator on
+accelerator cuda blah
+
+Description:
+
+Alter settings for use of accelerated versions of various styles.
+LAMMPS can be built with optional packages which provide accelerated
+versions of specific atom, pair,
+fix, compute, and integrate
+styles.
+
+These are the relevant packages:
+
+- OPT = a handful of pair styles, cache-optimized for faster CPU performance
+
- GPU = a handful of pair styles and the PPPM kspace_style, optimized to run on one or more GPUs or multicore CPU/GPU nodes
+
- USER-CUDA = a collection of atom, pair, fix, compute, and intergrate styles, optimized to run on one or more NVIDIA GPUs
+
+See this section of the manual for
+instructions on how to build LAMMPS with any of these packages.
+
+These styles can be specified explicitly in your input script,
+e.g. pair_style lj/cut/gpu. If the -accelerator command-line switch
+is used, you do not need to modify your input script, as discussed in
+this section of the manual. The command-line
+suffix (opt,gpu,cuda) is automatically appended when the style is
+created for atom, pair, fix, compute, and integrate styles. If an
+accelerated version does not exist, the standard version is created.
+
+If the -accelerator command-line switch is used, you may wish to
+disable it for one or more input script commands, so that the
+accelerated version of the style is not used, but rather the standard
+version. This can be useful for performance testing or debugging.
+
+The off style allows you to do this. The effect of the -accelerator
+command-line switch is effectively turned off until another
+accelerator command with the on style is used. The on style can
+only be used if the -accelerator command-line switch was used.
+
+The cuda style invokes options associated with the use of the
+USER-CUDA package. These will be described when the USER-CUDA package
+is released with LAMMPS.
+
+Restrictions:
+
+This cuda style can only be invoked if LAMMPS was built with the
+USER-CUDA package. See the Making LAMMPS
+section for more info.
+
+Obviously, you must have GPU hardware and associated software to
+build LAMMPS with GPU support.
+
+Related commands:
+
+fix gpu
+
+Default: none
+
+
diff --git a/doc/accelerator.txt b/doc/accelerator.txt
new file mode 100644
index 0000000000..c71f57c0a4
--- /dev/null
+++ b/doc/accelerator.txt
@@ -0,0 +1,80 @@
+"LAMMPS WWW Site"_lws - "LAMMPS Documentation"_ld - "LAMMPS Commands"_lc :c
+
+:link(lws,http://lammps.sandia.gov)
+:link(ld,Manual.html)
+:link(lc,Section_commands.html#comm)
+
+:line
+
+accelerator command :h3
+
+[Syntax:]
+
+accelerator style args :pre
+
+style = {off} or {on} or {cuda} :ulb,l
+args = 0 or more args specific to the style :l
+ {off} args = none
+ {on} args = none
+ {cuda} args = to be determined :pre
+:ule
+
+[Examples:]
+
+accelerator off
+accelerator on
+accelerator cuda blah :pre
+
+[Description:]
+
+Alter settings for use of accelerated versions of various styles.
+LAMMPS can be built with optional packages which provide accelerated
+versions of specific "atom"_atom_style.html, "pair"_pair_style.html,
+"fix"_fix.html, "compute"_compute.html, and "integrate"_run_style.html
+styles.
+
+These are the relevant packages:
+
+OPT = a handful of pair styles, cache-optimized for faster CPU performance
+GPU = a handful of pair styles and the PPPM kspace_style, optimized to run on one or more GPUs or multicore CPU/GPU nodes
+USER-CUDA = a collection of atom, pair, fix, compute, and intergrate styles, optimized to run on one or more NVIDIA GPUs :ul
+
+See "this section"_Section_start.html#2_3 of the manual for
+instructions on how to build LAMMPS with any of these packages.
+
+These styles can be specified explicitly in your input script,
+e.g. pair_style lj/cut/gpu. If the -accelerator command-line switch
+is used, you do not need to modify your input script, as discussed in
+"this section"_Section_start.html#2_6 of the manual. The command-line
+suffix (opt,gpu,cuda) is automatically appended when the style is
+created for atom, pair, fix, compute, and integrate styles. If an
+accelerated version does not exist, the standard version is created.
+
+If the -accelerator command-line switch is used, you may wish to
+disable it for one or more input script commands, so that the
+accelerated version of the style is not used, but rather the standard
+version. This can be useful for performance testing or debugging.
+
+The {off} style allows you to do this. The effect of the -accelerator
+command-line switch is effectively turned off until another
+accelerator command with the {on} style is used. The {on} style can
+only be used if the -accelerator command-line switch was used.
+
+The {cuda} style invokes options associated with the use of the
+USER-CUDA package. These will be described when the USER-CUDA package
+is released with LAMMPS.
+
+[Restrictions:]
+
+This cuda style can only be invoked if LAMMPS was built with the
+USER-CUDA package. See the "Making LAMMPS"_Section_start.html#2_3
+section for more info.
+
+Obviously, you must have GPU hardware and associated software to
+build LAMMPS with GPU support.
+
+[Related commands:]
+
+"fix gpu"_fix_gpu.html
+
+[Default:] none
From 813ffdc80dcffed099655f3274b752c277fddf1b Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 23:25:17 +0000
Subject: [PATCH 71/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6198
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/Section_start.html | 28 +++++++++++++++-------------
doc/Section_start.txt | 28 +++++++++++++++-------------
doc/accelerator.html | 21 +++++++++++----------
doc/accelerator.txt | 21 +++++++++++----------
4 files changed, 52 insertions(+), 46 deletions(-)
diff --git a/doc/Section_start.html b/doc/Section_start.html
index 999bb8b732..55e6bb352d 100644
--- a/doc/Section_start.html
+++ b/doc/Section_start.html
@@ -806,19 +806,21 @@ one-letter abbreviation can be used:
Use accelerated variants of various styles if they exist. The style
can be opt or gpu or cuda. The variant styles are part of
optional packages that LAMMPS can be built with, as described above in
-Section 2.3. The "opt" style corrsponds to the OPT package,
-the "gpu" style to the GPU package, and the "cuda" style to the
-USER-CUDA package. For example, all of the packages provide a
-pair_style lj/cut variant, with styles lj/cut/opt or
-lj/cut/gpu or lj/cut/cuda. These styles can be specified explicitly
-in your input script, e.g. pair_style lj/cut/gpu. But if the
--accelerator switch is used, you do not need to modify your input
-script. The accelerator suffix (opt,gpu,cuda) is automatically
-appended when the style is created. The suffix is applied to atom,
-pair, fix, compute, and integrate styles. If an accelerated version
-does not exist, the standard version is created. See the
-accelerator command for info on how to turn off
-this option.
+Section 2.3. Also see the acclerator
+command doc page. The "opt" style corrsponds to the OPT package, the
+"gpu" style to the GPU package, and the "cuda" style to the USER-CUDA
+package. For example, all of the packages provide a pair_style
+lj/cut variant, with style names lj/cut/opt or
+lj/cut/gpu or lj/cut/cuda.
+
+These accelerated styles can be specified explicitly in your input
+script, e.g. pair_style lj/cut/gpu. If the -accelerator switch is
+used, you do not need to modify your input script. The accelerator
+suffix (opt,gpu,cuda) is automatically appended when the style is
+created for atom, pair, fix, compute, and integrate styles. If an
+accelerated version does not exist, the standard version is created.
+See the accelerator command for info on how to
+temporarily turn off this option.
-echo style
diff --git a/doc/Section_start.txt b/doc/Section_start.txt
index 57bb2a8975..ea0462ae9a 100644
--- a/doc/Section_start.txt
+++ b/doc/Section_start.txt
@@ -796,19 +796,21 @@ Here are the details on the options:
Use accelerated variants of various styles if they exist. The style
can be {opt} or {gpu} or {cuda}. The variant styles are part of
optional packages that LAMMPS can be built with, as described above in
-"Section 2.3"_#2_3. The "opt" style corrsponds to the OPT package,
-the "gpu" style to the GPU package, and the "cuda" style to the
-USER-CUDA package. For example, all of the packages provide a
-"pair_style lj/cut"_pair_lj.html variant, with styles lj/cut/opt or
-lj/cut/gpu or lj/cut/cuda. These styles can be specified explicitly
-in your input script, e.g. pair_style lj/cut/gpu. But if the
--accelerator switch is used, you do not need to modify your input
-script. The accelerator suffix (opt,gpu,cuda) is automatically
-appended when the style is created. The suffix is applied to atom,
-pair, fix, compute, and integrate styles. If an accelerated version
-does not exist, the standard version is created. See the
-"accelerator"_accelerator.html command for info on how to turn off
-this option.
+"Section 2.3"_#2_3. Also see the "acclerator"_accelerator.html
+command doc page. The "opt" style corrsponds to the OPT package, the
+"gpu" style to the GPU package, and the "cuda" style to the USER-CUDA
+package. For example, all of the packages provide a "pair_style
+lj/cut"_pair_lj.html variant, with style names lj/cut/opt or
+lj/cut/gpu or lj/cut/cuda.
+
+These accelerated styles can be specified explicitly in your input
+script, e.g. pair_style lj/cut/gpu. If the -accelerator switch is
+used, you do not need to modify your input script. The accelerator
+suffix (opt,gpu,cuda) is automatically appended when the style is
+created for atom, pair, fix, compute, and integrate styles. If an
+accelerated version does not exist, the standard version is created.
+See the "accelerator"_accelerator.html command for info on how to
+temporarily turn off this option.
-echo style :pre
diff --git a/doc/accelerator.html b/doc/accelerator.html
index d53813863f..ff6eca3ec0 100644
--- a/doc/accelerator.html
+++ b/doc/accelerator.html
@@ -49,21 +49,22 @@ styles.
instructions on how to build LAMMPS with any of these packages.
These styles can be specified explicitly in your input script,
-e.g. pair_style lj/cut/gpu. If the -accelerator command-line switch
-is used, you do not need to modify your input script, as discussed in
-this section of the manual. The command-line
-suffix (opt,gpu,cuda) is automatically appended when the style is
-created for atom, pair, fix, compute, and integrate styles. If an
-accelerated version does not exist, the standard version is created.
+e.g. pair_style lj/cut/gpu. If the -accelerator
+command-line switch is used, you do not need to modify your input
+script, as discussed in this section of the
+manual. The command-line suffix (opt,gpu,cuda) is automatically
+appended when the style is created for atom, pair, fix, compute, and
+integrate styles. If an accelerated version does not exist, the
+standard version is created.
If the -accelerator command-line switch is used, you may wish to
-disable it for one or more input script commands, so that the
-accelerated version of the style is not used, but rather the standard
-version. This can be useful for performance testing or debugging.
+disable it for one or more input script commands, so that the standard
+version of the style is used instead of the accelerated one. This can
+be useful for performance testing or debugging.
The off style allows you to do this. The effect of the -accelerator
command-line switch is effectively turned off until another
-accelerator command with the on style is used. The on style can
+accelerator command is used with the on style. The on style can
only be used if the -accelerator command-line switch was used.
The cuda style invokes options associated with the use of the
diff --git a/doc/accelerator.txt b/doc/accelerator.txt
index c71f57c0a4..6f93f33a45 100644
--- a/doc/accelerator.txt
+++ b/doc/accelerator.txt
@@ -43,21 +43,22 @@ See "this section"_Section_start.html#2_3 of the manual for
instructions on how to build LAMMPS with any of these packages.
These styles can be specified explicitly in your input script,
-e.g. pair_style lj/cut/gpu. If the -accelerator command-line switch
-is used, you do not need to modify your input script, as discussed in
-"this section"_Section_start.html#2_6 of the manual. The command-line
-suffix (opt,gpu,cuda) is automatically appended when the style is
-created for atom, pair, fix, compute, and integrate styles. If an
-accelerated version does not exist, the standard version is created.
+e.g. "pair_style lj/cut/gpu"_pair_lj.html. If the -accelerator
+command-line switch is used, you do not need to modify your input
+script, as discussed in "this section"_Section_start.html#2_6 of the
+manual. The command-line suffix (opt,gpu,cuda) is automatically
+appended when the style is created for atom, pair, fix, compute, and
+integrate styles. If an accelerated version does not exist, the
+standard version is created.
If the -accelerator command-line switch is used, you may wish to
-disable it for one or more input script commands, so that the
-accelerated version of the style is not used, but rather the standard
-version. This can be useful for performance testing or debugging.
+disable it for one or more input script commands, so that the standard
+version of the style is used instead of the accelerated one. This can
+be useful for performance testing or debugging.
The {off} style allows you to do this. The effect of the -accelerator
command-line switch is effectively turned off until another
-accelerator command with the {on} style is used. The {on} style can
+accelerator command is used with the {on} style. The {on} style can
only be used if the -accelerator command-line switch was used.
The {cuda} style invokes options associated with the use of the
From 47eae4052fa99cdf3f65d9ce7eb249de507f9d7f Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 23:30:21 +0000
Subject: [PATCH 72/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6199
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index 110bee8220..f4b18f32dc 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "19 May 2011"
+#define LAMMPS_VERSION "20 May 2011"
From ad89319d7c5e11c491b766a037d54e2dcf3c0231 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 23:39:34 +0000
Subject: [PATCH 73/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6201
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/input.cpp | 7 ++++++-
1 file changed, 6 insertions(+), 1 deletion(-)
diff --git a/src/input.cpp b/src/input.cpp
index 6d38318089..370f6ceb8b 100644
--- a/src/input.cpp
+++ b/src/input.cpp
@@ -817,7 +817,12 @@ void Input::accelerator()
return;
}
- lmp->offaccel = 0;
+ if (strcmp(arg[0],"on") == 0) {
+ if (narg != 1) error->all("Illegal accelerator command");
+ lmp->offaccel = 0;
+ return;
+ }
+
if (strcmp(arg[0],"cuda") == 0) lmp->cuda->accelerator(narg-1,&arg[1]);
else error->all("Illegal accelerator command");
}
From 66c5a33eff54f1eeffa439fb3d1b266bb9a7a003 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Fri, 20 May 2011 23:47:42 +0000
Subject: [PATCH 74/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6202
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index f4b18f32dc..af815a44ba 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "20 May 2011"
+#define LAMMPS_VERSION "21 May 2011"
From c6e661ddba3e71788a963cba49ddaab52c7550aa Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Sat, 21 May 2011 00:30:37 +0000
Subject: [PATCH 75/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6204
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/ASPHERE/Install.sh | 11 -----------
src/Depend.sh | 30 ++++++++++++++++++++++++++++++
src/GPU/Install.sh | 23 +++++++++++------------
src/KSPACE/Install.sh | 7 -------
src/MANYBODY/Install.sh | 11 -----------
src/MEAM/Install.sh | 2 +-
src/Makefile | 4 ++--
src/OPT/Install.sh | 2 +-
src/POEMS/Install.sh | 2 +-
src/REAX/Install.sh | 2 +-
src/USER-ATC/Install.sh | 2 +-
src/USER-CG-CMM/Install.sh | 19 ++++++++++---------
12 files changed, 58 insertions(+), 57 deletions(-)
create mode 100644 src/Depend.sh
diff --git a/src/ASPHERE/Install.sh b/src/ASPHERE/Install.sh
index 7de96460f3..3faf9fe241 100644
--- a/src/ASPHERE/Install.sh
+++ b/src/ASPHERE/Install.sh
@@ -1,7 +1,4 @@
# Install/unInstall package files in LAMMPS
-# for unInstall, also unInstall/Install GPU package if installed
-# so it will remove GPU files that depend on ASPHERE files,
-# then replace others
if (test $1 = 1) then
@@ -25,10 +22,6 @@ if (test $1 = 1) then
cp pair_gayberne.h ..
cp pair_resquared.h ..
- if (test -e ../pair_lj_cut_gpu.h) then
- cd ../GPU; /bin/sh Install.sh 1
- fi
-
elif (test $1 = 0) then
rm ../compute_erotate_asphere.cpp
@@ -51,8 +44,4 @@ elif (test $1 = 0) then
rm ../pair_gayberne.h
rm ../pair_resquared.h
- if (test -e ../pair_gayberne_gpu.h) then
- cd ../GPU; /bin/sh Install.sh 0; /bin/sh Install.sh 1
- fi
-
fi
diff --git a/src/Depend.sh b/src/Depend.sh
new file mode 100644
index 0000000000..820b64070c
--- /dev/null
+++ b/src/Depend.sh
@@ -0,0 +1,30 @@
+# Depend.sh = Install/unInstall dependent package files
+# only Install/unInstall if dependent package is already installed
+# install dependent child files when new parent files installed
+# uninstall dependent child files when parent files uninstalled
+
+if (test $1 = 1) then
+
+ if (test -e pair_lj_cut_opt.h) then
+ cd OPT; /bin/sh Install.sh 1; cd ..
+ fi
+ if (test -e pair_lj_cut_gpu.h) then
+ cd GPU; /bin/sh Install.sh 1; cd ..
+ fi
+ if (test -e pair_lj_cut_cuda.h) then
+ cd USER-CUDA; /bin/sh Install.sh 1; cd ..
+ fi
+
+elif (test $1 = 0) then
+
+ if (test -e pair_lj_cut_opt.h) then
+ cd OPT; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
+ fi
+ if (test -e pair_lj_cut_gpu.h) then
+ cd GPU; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
+ fi
+ if (test -e pair_lj_cut_cuda.h) then
+ cd USER-CUDA; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
+ fi
+
+fi
diff --git a/src/GPU/Install.sh b/src/GPU/Install.sh
index 6367a6cea5..89dfa6b0cb 100644
--- a/src/GPU/Install.sh
+++ b/src/GPU/Install.sh
@@ -1,7 +1,6 @@
# Install/unInstall package files in LAMMPS
-# edit Makefile.package to include/exclude GPU library
-# do not copy gayberne files if non-GPU version does not exist
-# do not copy charmm files if non-GPU version does not exist
+# edit Makefile.package to include/exclude GPU info
+# do not install child files if parent does not exist
if (test $1 = 1) then
@@ -15,15 +14,6 @@ if (test $1 = 1) then
sed -i -e 's|^PKG_SYSPATH =[ \t]*|&$(gpu_SYSPATH) |' ../Makefile.package
fi
- if (test -e ../pppm.cpp) then
- cp pppm_gpu.cpp ..
- cp pppm_gpu_single.cpp ..
- cp pppm_gpu_double.cpp ..
- cp pppm_gpu.h ..
- cp pppm_gpu_single.h ..
- cp pppm_gpu_double.h ..
- fi
-
if (test -e ../pair_gayberne.cpp) then
cp pair_gayberne_gpu.cpp ..
cp pair_gayberne_gpu.h ..
@@ -49,6 +39,15 @@ if (test $1 = 1) then
cp pair_cg_cmm_coul_long_gpu.h ..
fi
+ if (test -e ../pppm.cpp) then
+ cp pppm_gpu.cpp ..
+ cp pppm_gpu_single.cpp ..
+ cp pppm_gpu_double.cpp ..
+ cp pppm_gpu.h ..
+ cp pppm_gpu_single.h ..
+ cp pppm_gpu_double.h ..
+ fi
+
cp pair_lj_cut_gpu.cpp ..
cp pair_morse_gpu.cpp ..
cp pair_lj96_cut_gpu.cpp ..
diff --git a/src/KSPACE/Install.sh b/src/KSPACE/Install.sh
index efecdf33f8..4db0f3e010 100644
--- a/src/KSPACE/Install.sh
+++ b/src/KSPACE/Install.sh
@@ -1,7 +1,4 @@
# Install/unInstall package files in LAMMPS
-# for unInstall, also unInstall/Install OPT package if installed
-# so it will remove OPT files that depend on KSPACE files,
-# then replace others
if (test $1 = 1) then
@@ -63,8 +60,4 @@ elif (test $1 = 0) then
rm ../remap.h
rm ../remap_wrap.h
- if (test -e ../pair_lj_charmm_coul_long_opt.h) then
- cd ../OPT; sh Install.sh 0; sh Install.sh 1
- fi
-
fi
diff --git a/src/MANYBODY/Install.sh b/src/MANYBODY/Install.sh
index 0817b08cb2..12a551a6f8 100644
--- a/src/MANYBODY/Install.sh
+++ b/src/MANYBODY/Install.sh
@@ -1,7 +1,4 @@
# Install/unInstall package files in LAMMPS
-# for unInstall, also unInstall/Install OPT package if installed
-# so it will remove OPT files that depend on MANYBODY files,
-# then replace others
if (test $1 = 1) then
@@ -27,10 +24,6 @@ if (test $1 = 1) then
cp pair_tersoff.h ..
cp pair_tersoff_zbl.h ..
- if (test -e ../pair_lj_cut_opt.h) then
- cd ../OPT; sh Install.sh 1
- fi
-
elif (test $1 = 0) then
rm ../fix_qeq_comb.cpp
@@ -55,8 +48,4 @@ elif (test $1 = 0) then
rm ../pair_tersoff.h
rm ../pair_tersoff_zbl.h
- if (test -e ../pair_eam_opt.h) then
- cd ../OPT; sh Install.sh 0; sh Install.sh 1
- fi
-
fi
diff --git a/src/MEAM/Install.sh b/src/MEAM/Install.sh
index 4ce114f3c4..a2ff3ad029 100644
--- a/src/MEAM/Install.sh
+++ b/src/MEAM/Install.sh
@@ -1,5 +1,5 @@
# Install/unInstall package files in LAMMPS
-# edit Makefile.package to include/exclude MEAM library
+# edit Makefile.package to include/exclude MEAM info
if (test $1 = 1) then
diff --git a/src/Makefile b/src/Makefile
index 7473b0e112..1812e979b9 100755
--- a/src/Makefile
+++ b/src/Makefile
@@ -149,7 +149,7 @@ yes-%:
echo "Package $(@:yes-%=%) does not exist"; \
else \
echo "Installing package $(@:yes-%=%)"; \
- cd $(YESDIR); $(SHELL) Install.sh 1; \
+ cd $(YESDIR); $(SHELL) Install.sh 1; cd ..; $(SHELL) Depend.sh 1; \
fi;
no-%:
@@ -157,7 +157,7 @@ no-%:
echo "Package $(@:no-%=%) does not exist"; \
else \
echo "Uninstalling package $(@:no-%=%), ignore errors"; \
- cd $(NODIR); $(SHELL) Install.sh 0; cd ..; \
+ cd $(NODIR); $(SHELL) Install.sh 0; cd ..; $(SHELL) Depend.sh 0; \
fi;
# status = list differences between src and package files
diff --git a/src/OPT/Install.sh b/src/OPT/Install.sh
index 4b12d0396a..37cedbb506 100644
--- a/src/OPT/Install.sh
+++ b/src/OPT/Install.sh
@@ -1,5 +1,5 @@
# Install/unInstall package files in LAMMPS
-# do not copy eam and charmm files if non-OPT versions do not exist
+# do not install child files if parent does not exist
if (test $1 = 1) then
diff --git a/src/POEMS/Install.sh b/src/POEMS/Install.sh
index 58a33048c7..59a86ee924 100644
--- a/src/POEMS/Install.sh
+++ b/src/POEMS/Install.sh
@@ -1,5 +1,5 @@
# Install/unInstall package files in LAMMPS
-# edit Makefile.package to include/exclude POEMS library
+# edit Makefile.package to include/exclude POEMS info
if (test $1 = 1) then
diff --git a/src/REAX/Install.sh b/src/REAX/Install.sh
index abf9078568..69c55f60c2 100644
--- a/src/REAX/Install.sh
+++ b/src/REAX/Install.sh
@@ -1,5 +1,5 @@
# Install/unInstall package files in LAMMPS
-# edit Makefile.package to include/exclude REAX library
+# edit Makefile.package to include/exclude REAX info
if (test $1 = 1) then
diff --git a/src/USER-ATC/Install.sh b/src/USER-ATC/Install.sh
index 307b9cab00..9bb438f6a0 100755
--- a/src/USER-ATC/Install.sh
+++ b/src/USER-ATC/Install.sh
@@ -1,5 +1,5 @@
# Install/unInstall package files in LAMMPS
-# edit Makefile.package to include/exclude ATC library
+# edit Makefile.package to include/exclude ATC info
if (test $1 = 1) then
diff --git a/src/USER-CG-CMM/Install.sh b/src/USER-CG-CMM/Install.sh
index 18dcb44d3a..ada215a923 100644
--- a/src/USER-CG-CMM/Install.sh
+++ b/src/USER-CG-CMM/Install.sh
@@ -1,4 +1,5 @@
# Install/unInstall package files in LAMMPS
+# do not install child files if parent does not exist
if (test $1 = 1) then
@@ -7,6 +8,11 @@ if (test $1 = 1) then
cp angle_cg_cmm.cpp ..
fi
+ if (test -e ../pppm.cpp) then
+ cp pair_cg_cmm_coul_long.cpp ..
+ cp pair_cg_cmm_coul_long.h ..
+ fi
+
cp cg_cmm_parms.h ..
cp cg_cmm_parms.cpp ..
@@ -17,15 +23,10 @@ if (test $1 = 1) then
cp pair_cg_cmm_coul_cut.cpp ..
cp pair_cg_cmm_coul_cut.h ..
- if (test -e ../pppm.cpp) then
- cp pair_cg_cmm_coul_long.cpp ..
- cp pair_cg_cmm_coul_long.h ..
- fi
-
elif (test $1 = 0) then
- rm -f ../angle_cg_cmm.h
- rm -f ../angle_cg_cmm.cpp
+ rm ../angle_cg_cmm.h
+ rm ../angle_cg_cmm.cpp
rm ../cg_cmm_parms.h
rm ../cg_cmm_parms.cpp
@@ -37,7 +38,7 @@ elif (test $1 = 0) then
rm ../pair_cg_cmm_coul_cut.cpp
rm ../pair_cg_cmm_coul_cut.h
- rm -f ../pair_cg_cmm_coul_long.cpp
- rm -f ../pair_cg_cmm_coul_long.h
+ rm ../pair_cg_cmm_coul_long.cpp
+ rm ../pair_cg_cmm_coul_long.h
fi
From de61f3db323c9c83fe930fa6bc6010e59b4ecb63 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Sat, 21 May 2011 00:38:51 +0000
Subject: [PATCH 76/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6205
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/Depend.sh | 10 ++++++++--
1 file changed, 8 insertions(+), 2 deletions(-)
diff --git a/src/Depend.sh b/src/Depend.sh
index 820b64070c..f6f3d43d08 100644
--- a/src/Depend.sh
+++ b/src/Depend.sh
@@ -1,6 +1,6 @@
-# Depend.sh = Install/unInstall dependent package files
+# Depend.sh = Install/unInstall files from dependent packages
# only Install/unInstall if dependent package is already installed
-# install dependent child files when new parent files installed
+# install dependent child files when parent files installed
# uninstall dependent child files when parent files uninstalled
if (test $1 = 1) then
@@ -11,6 +11,9 @@ if (test $1 = 1) then
if (test -e pair_lj_cut_gpu.h) then
cd GPU; /bin/sh Install.sh 1; cd ..
fi
+ if (test -e cg_cmm_params.h) then
+ cd USER-CG-CMM; /bin/sh Install.sh 1; cd ..
+ fi
if (test -e pair_lj_cut_cuda.h) then
cd USER-CUDA; /bin/sh Install.sh 1; cd ..
fi
@@ -23,6 +26,9 @@ elif (test $1 = 0) then
if (test -e pair_lj_cut_gpu.h) then
cd GPU; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
fi
+ if (test -e cg_cmm_params.h) then
+ cd USER-CG-CMM; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
+ fi
if (test -e pair_lj_cut_cuda.h) then
cd USER-CUDA; /bin/sh Install.sh 0; /bin/sh Install.sh 1; cd ..
fi
From 415a308210ecd027ac04f812347320f1658f3a73 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Sat, 21 May 2011 00:40:41 +0000
Subject: [PATCH 77/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6206
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index af815a44ba..49d208702f 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "21 May 2011"
+#define LAMMPS_VERSION "22 May 2011"
From 69786dc452eb69abd3448eca72fa7e5fa2079017 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Mon, 23 May 2011 13:21:38 +0000
Subject: [PATCH 78/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6208
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/accelerator.h | 73 +++++++++++++++++++++++++++++++++++++++++++++++
1 file changed, 73 insertions(+)
create mode 100644 src/accelerator.h
diff --git a/src/accelerator.h b/src/accelerator.h
new file mode 100644
index 0000000000..9ef291adbc
--- /dev/null
+++ b/src/accelerator.h
@@ -0,0 +1,73 @@
+/* ----------------------------------------------------------------------
+ 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.
+------------------------------------------------------------------------- */
+
+// dummy interface to USER-CUDA
+// used when USER-CUDA is not installed
+
+#ifndef LMP_ACCELERATOR_H
+#define LMP_ACCELERATOR_H
+
+#include "comm.h"
+#include "modify.h"
+#include "verlet.h"
+
+namespace LAMMPS_NS {
+
+class Cuda {
+ public:
+ int cuda_exists;
+ int oncpu;
+ int neighbor_decide_by_integrator;
+
+ Cuda(class LAMMPS *) {cuda_exists = 0;}
+ ~Cuda() {}
+ void setDevice(class LAMMPS *) {}
+ void accelerator(int, char **) {}
+ void evsetup_eatom_vatom(int, int) {}
+ void downloadAll() {}
+ void uploadAll() {}
+};
+
+class CommCuda : public Comm {
+ public:
+ CommCuda(class LAMMPS *lmp) : Comm(lmp) {}
+ ~CommCuda() {}
+};
+
+class DomainCuda : public Domain {
+ public:
+ DomainCuda(class LAMMPS *lmp) : Domain(lmp) {}
+ ~DomainCuda() {}
+};
+
+class NeighborCuda : public Neighbor {
+ public:
+ NeighborCuda(class LAMMPS *lmp) : Neighbor(lmp) {}
+ ~NeighborCuda() {}
+};
+
+class ModifyCuda : public Modify {
+ public:
+ ModifyCuda(class LAMMPS *lmp) : Modify(lmp) {}
+ ~ModifyCuda() {}
+};
+
+class VerletCuda : public Verlet {
+ public:
+ VerletCuda(class LAMMPS *lmp, int narg, char **arg) : Verlet(lmp,narg,arg) {}
+ ~VerletCuda() {}
+};
+
+}
+
+#endif
From e97e40698a19fe067d8946fc41ae25097ee249eb Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Mon, 23 May 2011 14:04:44 +0000
Subject: [PATCH 79/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6209
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/MOLECULE/atom_vec_angle.h | 22 +++++++--------
src/MOLECULE/atom_vec_full.h | 22 +++++++--------
src/atom.cpp | 4 +--
src/atom_vec_atomic.h | 22 +++++++--------
src/atom_vec_charge.h | 22 +++++++--------
src/comm.h | 50 +++++++++++++++++------------------
src/domain.h | 18 ++++++-------
src/input.cpp | 2 +-
src/lammps.cpp | 3 ++-
src/modify.h | 28 ++++++++++----------
src/output.cpp | 13 +++++++++
src/update.cpp | 16 ++++++++---
12 files changed, 122 insertions(+), 100 deletions(-)
diff --git a/src/MOLECULE/atom_vec_angle.h b/src/MOLECULE/atom_vec_angle.h
index a9f35573ee..cfe83c716e 100644
--- a/src/MOLECULE/atom_vec_angle.h
+++ b/src/MOLECULE/atom_vec_angle.h
@@ -27,24 +27,24 @@ namespace LAMMPS_NS {
class AtomVecAngle : public AtomVec {
public:
AtomVecAngle(class LAMMPS *, int, char **);
- virtual ~AtomVecAngle() {}
+ ~AtomVecAngle() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
- int pack_comm(int, int *, double *, int, int *);
- int pack_comm_vel(int, int *, double *, int, int *);
- void unpack_comm(int, int, double *);
- void unpack_comm_vel(int, int, double *);
+ virtual int pack_comm(int, int *, double *, int, int *);
+ virtual int pack_comm_vel(int, int *, double *, int, int *);
+ virtual void unpack_comm(int, int, double *);
+ virtual void unpack_comm_vel(int, int, double *);
int pack_reverse(int, int, double *);
void unpack_reverse(int, int *, double *);
- int pack_border(int, int *, double *, int, int *);
- int pack_border_vel(int, int *, double *, int, int *);
+ virtual int pack_border(int, int *, double *, int, int *);
+ virtual int pack_border_vel(int, int *, double *, int, int *);
int pack_border_hybrid(int, int *, double *);
- void unpack_border(int, int, double *);
- void unpack_border_vel(int, int, double *);
+ virtual void unpack_border(int, int, double *);
+ virtual void unpack_border_vel(int, int, double *);
int unpack_border_hybrid(int, int, double *);
- int pack_exchange(int, double *);
- int unpack_exchange(double *);
+ virtual int pack_exchange(int, double *);
+ virtual int unpack_exchange(double *);
int size_restart();
int pack_restart(int, double *);
int unpack_restart(double *);
diff --git a/src/MOLECULE/atom_vec_full.h b/src/MOLECULE/atom_vec_full.h
index de68deb86c..a457c04660 100644
--- a/src/MOLECULE/atom_vec_full.h
+++ b/src/MOLECULE/atom_vec_full.h
@@ -27,24 +27,24 @@ namespace LAMMPS_NS {
class AtomVecFull : public AtomVec {
public:
AtomVecFull(class LAMMPS *, int, char **);
- virtual ~AtomVecFull() {}
+ ~AtomVecFull() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
- int pack_comm(int, int *, double *, int, int *);
- int pack_comm_vel(int, int *, double *, int, int *);
- void unpack_comm(int, int, double *);
- void unpack_comm_vel(int, int, double *);
+ virtual int pack_comm(int, int *, double *, int, int *);
+ virtual int pack_comm_vel(int, int *, double *, int, int *);
+ virtual void unpack_comm(int, int, double *);
+ virtual void unpack_comm_vel(int, int, double *);
int pack_reverse(int, int, double *);
void unpack_reverse(int, int *, double *);
- int pack_border(int, int *, double *, int, int *);
- int pack_border_vel(int, int *, double *, int, int *);
+ virtual int pack_border(int, int *, double *, int, int *);
+ virtual int pack_border_vel(int, int *, double *, int, int *);
int pack_border_hybrid(int, int *, double *);
- void unpack_border(int, int, double *);
- void unpack_border_vel(int, int, double *);
+ virtual void unpack_border(int, int, double *);
+ virtual void unpack_border_vel(int, int, double *);
int unpack_border_hybrid(int, int, double *);
- int pack_exchange(int, double *);
- int unpack_exchange(double *);
+ virtual int pack_exchange(int, double *);
+ virtual int unpack_exchange(double *);
int size_restart();
int pack_restart(int, double *);
int unpack_restart(double *);
diff --git a/src/atom.cpp b/src/atom.cpp
index 9abc972f3c..d7c08d9cbb 100644
--- a/src/atom.cpp
+++ b/src/atom.cpp
@@ -1446,11 +1446,11 @@ void Atom::setup_sort_bins()
double vol = (domain->boxhi[0]-domain->boxlo[0]) *
(domain->boxhi[1]-domain->boxlo[1]) *
(domain->boxhi[2]-domain->boxlo[2]);
- binsize = pow(CUDA_CHUNK/natoms*vol,1.0/3.0);
+ binsize = pow(1.0*CUDA_CHUNK/natoms*vol,1.0/3.0);
} else {
double area = (domain->boxhi[0]-domain->boxlo[0]) *
(domain->boxhi[1]-domain->boxlo[1]);
- binsize = pow(CUDA_CHUNK/natoms*area,1.0/2.0);
+ binsize = pow(1.0*CUDA_CHUNK/natoms*area,1.0/2.0);
}
} else binsize = 0.5 * neighbor->cutneighmax;
if (binsize == 0.0) error->all("Atom sorting has bin size = 0.0");
diff --git a/src/atom_vec_atomic.h b/src/atom_vec_atomic.h
index 177038729b..537065c553 100644
--- a/src/atom_vec_atomic.h
+++ b/src/atom_vec_atomic.h
@@ -27,22 +27,22 @@ namespace LAMMPS_NS {
class AtomVecAtomic : public AtomVec {
public:
AtomVecAtomic(class LAMMPS *, int, char **);
- virtual ~AtomVecAtomic() {}
+ ~AtomVecAtomic() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
- int pack_comm(int, int *, double *, int, int *);
- int pack_comm_vel(int, int *, double *, int, int *);
- void unpack_comm(int, int, double *);
- void unpack_comm_vel(int, int, double *);
+ virtual int pack_comm(int, int *, double *, int, int *);
+ virtual int pack_comm_vel(int, int *, double *, int, int *);
+ virtual void unpack_comm(int, int, double *);
+ virtual void unpack_comm_vel(int, int, double *);
int pack_reverse(int, int, double *);
void unpack_reverse(int, int *, double *);
- int pack_border(int, int *, double *, int, int *);
- int pack_border_vel(int, int *, double *, int, int *);
- void unpack_border(int, int, double *);
- void unpack_border_vel(int, int, double *);
- int pack_exchange(int, double *);
- int unpack_exchange(double *);
+ virtual int pack_border(int, int *, double *, int, int *);
+ virtual int pack_border_vel(int, int *, double *, int, int *);
+ virtual void unpack_border(int, int, double *);
+ virtual void unpack_border_vel(int, int, double *);
+ virtual int pack_exchange(int, double *);
+ virtual int unpack_exchange(double *);
int size_restart();
int pack_restart(int, double *);
int unpack_restart(double *);
diff --git a/src/atom_vec_charge.h b/src/atom_vec_charge.h
index cc37810598..cfd6555acb 100644
--- a/src/atom_vec_charge.h
+++ b/src/atom_vec_charge.h
@@ -27,24 +27,24 @@ namespace LAMMPS_NS {
class AtomVecCharge : public AtomVec {
public:
AtomVecCharge(class LAMMPS *, int, char **);
- virtual ~AtomVecCharge() {}
+ ~AtomVecCharge() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
- int pack_comm(int, int *, double *, int, int *);
- int pack_comm_vel(int, int *, double *, int, int *);
- void unpack_comm(int, int, double *);
- void unpack_comm_vel(int, int, double *);
+ virtual int pack_comm(int, int *, double *, int, int *);
+ virtual int pack_comm_vel(int, int *, double *, int, int *);
+ virtual void unpack_comm(int, int, double *);
+ virtual void unpack_comm_vel(int, int, double *);
int pack_reverse(int, int, double *);
void unpack_reverse(int, int *, double *);
- int pack_border(int, int *, double *, int, int *);
- int pack_border_vel(int, int *, double *, int, int *);
+ virtual int pack_border(int, int *, double *, int, int *);
+ virtual int pack_border_vel(int, int *, double *, int, int *);
int pack_border_hybrid(int, int *, double *);
- void unpack_border(int, int, double *);
- void unpack_border_vel(int, int, double *);
+ virtual void unpack_border(int, int, double *);
+ virtual void unpack_border_vel(int, int, double *);
int unpack_border_hybrid(int, int, double *);
- int pack_exchange(int, double *);
- int unpack_exchange(double *);
+ virtual int pack_exchange(int, double *);
+ virtual int unpack_exchange(double *);
int size_restart();
int pack_restart(int, double *);
int unpack_restart(double *);
diff --git a/src/comm.h b/src/comm.h
index d6934775bf..868f6bec2e 100644
--- a/src/comm.h
+++ b/src/comm.h
@@ -37,23 +37,23 @@ class Comm : protected Pointers {
Comm(class LAMMPS *);
virtual ~Comm();
- void init();
- void set_procs(); // setup 3d grid of procs
- void setup(); // setup 3d communication pattern
- void forward_comm(int dummy = 0); // forward communication of atom coords
- void reverse_comm(); // reverse communication of forces
- void exchange(); // move atoms to new procs
- void borders(); // setup list of atoms to communicate
+ virtual void init();
+ virtual void set_procs(); // setup 3d grid of procs
+ virtual void setup(); // setup 3d communication pattern
+ virtual void forward_comm(int dummy = 0); // forward communication of atom coords
+ virtual void reverse_comm(); // reverse communication of forces
+ virtual void exchange(); // move atoms to new procs
+ virtual void borders(); // setup list of atoms to communicate
- void forward_comm_pair(class Pair *); // forward comm from a Pair
- void reverse_comm_pair(class Pair *); // reverse comm from a Pair
- void forward_comm_fix(class Fix *); // forward comm from a Fix
- void reverse_comm_fix(class Fix *); // reverse comm from a Fix
- void forward_comm_compute(class Compute *); // forward comm from a Compute
- void reverse_comm_compute(class Compute *); // reverse comm from a Compute
+ virtual void forward_comm_pair(class Pair *); // forward comm from a Pair
+ virtual void reverse_comm_pair(class Pair *); // reverse comm from a Pair
+ virtual void forward_comm_fix(class Fix *); // forward comm from a Fix
+ virtual void reverse_comm_fix(class Fix *); // reverse comm from a Fix
+ virtual void forward_comm_compute(class Compute *); // forward comm from a Compute
+ virtual void reverse_comm_compute(class Compute *); // reverse comm from a Compute
- void set(int, char **); // set communication style
- bigint memory_usage();
+ virtual void set(int, char **); // set communication style
+ virtual bigint memory_usage();
protected:
int style; // single vs multi-type comm
@@ -87,18 +87,18 @@ class Comm : protected Pointers {
int maxsend,maxrecv; // current size of send/recv buffer
int maxforward,maxreverse; // max # of datums in forward/reverse comm
- void procs2box(); // map procs to 3d box
- void cross(double, double, double,
+ virtual void procs2box(); // map procs to 3d box
+ virtual void cross(double, double, double,
double, double, double,
double &, double &, double &); // cross product
- void grow_send(int,int); // reallocate send buffer
- void grow_recv(int); // free/allocate recv buffer
- void grow_list(int, int); // reallocate one sendlist
- void grow_swap(int); // grow swap and multi arrays
- void allocate_swap(int); // allocate swap arrays
- void allocate_multi(int); // allocate multi arrays
- void free_swap(); // free swap arrays
- void free_multi(); // free multi arrays
+ virtual void grow_send(int,int); // reallocate send buffer
+ virtual void grow_recv(int); // free/allocate recv buffer
+ virtual void grow_list(int, int); // reallocate one sendlist
+ virtual void grow_swap(int); // grow swap and multi arrays
+ virtual void allocate_swap(int); // allocate swap arrays
+ virtual void allocate_multi(int); // allocate multi arrays
+ virtual void free_swap(); // free swap arrays
+ virtual void free_multi(); // free multi arrays
};
}
diff --git a/src/domain.h b/src/domain.h
index d7cb88ce5c..3e8eba5d3e 100644
--- a/src/domain.h
+++ b/src/domain.h
@@ -85,14 +85,14 @@ class Domain : protected Pointers {
class Region **regions; // list of defined Regions
Domain(class LAMMPS *);
- ~Domain();
- void init();
+ virtual ~Domain();
+ virtual void init();
void set_initial_box();
- void set_global_box();
- void set_lamda_box();
- void set_local_box();
- void reset_box();
- void pbc();
+ virtual void set_global_box();
+ virtual void set_lamda_box();
+ virtual void set_local_box();
+ virtual void reset_box();
+ virtual void pbc();
void remap(double *, int &);
void remap(double *);
void remap_near(double *, double *);
@@ -107,8 +107,8 @@ class Domain : protected Pointers {
void set_boundary(int, char **);
void print_box(const char *);
- void lamda2x(int);
- void x2lamda(int);
+ virtual void lamda2x(int);
+ virtual void x2lamda(int);
void lamda2x(double *, double *);
void x2lamda(double *, double *);
void bbox(double *, double *, double *, double *);
diff --git a/src/input.cpp b/src/input.cpp
index 370f6ceb8b..079879574f 100644
--- a/src/input.cpp
+++ b/src/input.cpp
@@ -808,7 +808,7 @@ void Input::accelerator()
if (domain->box_exist)
error->all("Accelerator command after simulation box is defined");
if (narg < 1) error->all("Illegal accelerator command");
- if (strcmp(lmp->asuffix,arg[0]) != 0)
+ if (!lmp->asuffix || (strcmp(lmp->asuffix,arg[0]) != 0))
error->all("Accelerator command requires matching command-line -a switch");
if (strcmp(arg[0],"off") == 0) {
diff --git a/src/lammps.cpp b/src/lammps.cpp
index daf9c5ac50..f0b79ec45b 100644
--- a/src/lammps.cpp
+++ b/src/lammps.cpp
@@ -107,6 +107,7 @@ LAMMPS::LAMMPS(int narg, char **arg, MPI_Comm communicator)
else error->universe_all("Invalid command-line argument");
asuffix = new char[8];
strcpy(asuffix,arg[iarg+1]);
+ iarg += 2;
} else error->universe_all("Invalid command-line argument");
}
@@ -363,7 +364,7 @@ void LAMMPS::create()
void LAMMPS::init()
{
- if (accelerator == USERCUDA) cuda->setDevice(this);
+ if (accelerator == USERCUDA) cuda->accelerator(0,NULL);
update->init();
force->init(); // pair must come after update due to minimizer
diff --git a/src/modify.h b/src/modify.h
index fdd25b1045..0175cc06a0 100644
--- a/src/modify.h
+++ b/src/modify.h
@@ -41,21 +41,21 @@ class Modify : protected Pointers {
Modify(class LAMMPS *);
virtual ~Modify();
- void init();
- void setup(int);
- void setup_pre_exchange();
- void setup_pre_force(int);
- void initial_integrate(int);
- void post_integrate();
+ virtual void init();
+ virtual void setup(int);
+ virtual void setup_pre_exchange();
+ virtual void setup_pre_force(int);
+ virtual void initial_integrate(int);
+ virtual void post_integrate();
void pre_decide();
- void pre_exchange();
- void pre_neighbor();
- void pre_force(int);
- void post_force(int);
- void final_integrate();
- void end_of_step();
- double thermo_energy();
- void post_run();
+ virtual void pre_exchange();
+ virtual void pre_neighbor();
+ virtual void pre_force(int);
+ virtual void post_force(int);
+ virtual void final_integrate();
+ virtual void end_of_step();
+ virtual double thermo_energy();
+ virtual void post_run();
void setup_pre_force_respa(int, int);
void initial_integrate_respa(int, int, int);
diff --git a/src/output.cpp b/src/output.cpp
index a7339c1f58..95808a9b5e 100644
--- a/src/output.cpp
+++ b/src/output.cpp
@@ -32,9 +32,12 @@
#include "write_restart.h"
#include "memory.h"
#include "error.h"
+#include "accelerator.h"
using namespace LAMMPS_NS;
+enum{NOACCEL,OPT,GPU,USERCUDA}; // same as lammps.cpp
+
#define DELTA 1
#define MYMIN(a,b) ((a) < (b) ? (a) : (b))
@@ -243,8 +246,13 @@ void Output::write(bigint ntimestep)
{
// next_dump does not force output on last step of run
// wrap dumps that invoke computes with clear/add
+ // download data from GPU if necessary
if (next_dump_any == ntimestep) {
+
+ if (lmp->accelerator == USERCUDA && !lmp->cuda->oncpu)
+ lmp->cuda->downloadAll();
+
for (int idump = 0; idump < ndump; idump++) {
if (next_dump[idump] == ntimestep && last_dump[idump] != ntimestep) {
if (dump[idump]->clearstep) modify->clearstep_compute();
@@ -267,8 +275,13 @@ void Output::write(bigint ntimestep)
// next_restart does not force output on last step of run
// for toggle = 0, replace "*" with current timestep in restart filename
+ // download data from GPU if necessary
if (next_restart == ntimestep && last_restart != ntimestep) {
+
+ if (lmp->accelerator == USERCUDA && !lmp->cuda->oncpu)
+ lmp->cuda->downloadAll();
+
if (restart_toggle == 0) {
char *file = new char[strlen(restart1) + 16];
char *ptr = strchr(restart1,'*');
diff --git a/src/update.cpp b/src/update.cpp
index 1c70ad7047..c8efe4b796 100644
--- a/src/update.cpp
+++ b/src/update.cpp
@@ -216,15 +216,20 @@ void Update::create_integrate(int narg, char **arg, char *suffix)
}
}
-/* ---------------------------------------------------------------------- */
+/* ----------------------------------------------------------------------
+ create the Integrate style, first with suffix appended
+------------------------------------------------------------------------- */
void Update::new_integrate(char *style, int narg, char **arg,
char *suffix, int &sflag)
{
+ int success = 0;
+
if (suffix && lmp->offaccel == 0) {
sflag = 1;
char estyle[256];
sprintf(estyle,"%s/%s",style,suffix);
+ success = 1;
if (0) return;
@@ -235,20 +240,23 @@ void Update::new_integrate(char *style, int narg, char **arg,
#undef IntegrateStyle
#undef INTEGRATE_CLASS
+ else success = 0;
}
sflag = 0;
- if (0) return;
+ if (!success) {
+ if (0) return;
#define INTEGRATE_CLASS
#define IntegrateStyle(key,Class) \
- else if (strcmp(style,#key) == 0) integrate = new Class(lmp,narg,arg);
+ else if (strcmp(style,#key) == 0) integrate = new Class(lmp,narg,arg);
#include "style_integrate.h"
#undef IntegrateStyle
#undef INTEGRATE_CLASS
- else error->all("Illegal integrate style");
+ else error->all("Illegal integrate style");
+ }
}
/* ---------------------------------------------------------------------- */
From 94884d513bcb2bf87908495207b36054265411cc Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Mon, 23 May 2011 14:06:11 +0000
Subject: [PATCH 80/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6210
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/version.h | 2 +-
1 file changed, 1 insertion(+), 1 deletion(-)
diff --git a/src/version.h b/src/version.h
index 49d208702f..2e56b98b95 100644
--- a/src/version.h
+++ b/src/version.h
@@ -1 +1 @@
-#define LAMMPS_VERSION "22 May 2011"
+#define LAMMPS_VERSION "23 May 2011"
From 1c79295aa1f38949fb0aa0e6a75efa972c5829bc Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Mon, 23 May 2011 14:35:19 +0000
Subject: [PATCH 81/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6212
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/compute_group_group.html | 2 +-
doc/compute_group_group.txt | 2 +-
2 files changed, 2 insertions(+), 2 deletions(-)
diff --git a/doc/compute_group_group.html b/doc/compute_group_group.html
index 57e867136a..bb665ca6f7 100644
--- a/doc/compute_group_group.html
+++ b/doc/compute_group_group.html
@@ -49,7 +49,7 @@ section for an overview of LAMMPS output
options.
Both the scalar and vector values calculated by this compute are
-"extensive"., The scalar value will be in energy units.
+"extensive". The scalar value will be in energy units.
The vector values will be in force units.
Restrictions:
diff --git a/doc/compute_group_group.txt b/doc/compute_group_group.txt
index 71bfcf6450..5a77bcd561 100644
--- a/doc/compute_group_group.txt
+++ b/doc/compute_group_group.txt
@@ -46,7 +46,7 @@ section"_Section_howto.html#4_15 for an overview of LAMMPS output
options.
Both the scalar and vector values calculated by this compute are
-"extensive"., The scalar value will be in energy "units"_units.html.
+"extensive". The scalar value will be in energy "units"_units.html.
The vector values will be in force "units"_units.html.
[Restrictions:]
From 2ddc39d3d8198eced11bc476cffe9c4d48848c89 Mon Sep 17 00:00:00 2001
From: athomps
Date: Mon, 23 May 2011 19:09:02 +0000
Subject: [PATCH 82/83] Added equations for inverse box transformation
git-svn-id: svn://svn.icms.temple.edu/lammps-ro/trunk@6214 f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
doc/Eqs/box_inverse.jpg | Bin 0 -> 4570 bytes
doc/Eqs/box_inverse.tex | 14 ++++++++++++++
doc/Section_howto.html | 8 ++++++--
3 files changed, 20 insertions(+), 2 deletions(-)
create mode 100644 doc/Eqs/box_inverse.jpg
create mode 100644 doc/Eqs/box_inverse.tex
diff --git a/doc/Eqs/box_inverse.jpg b/doc/Eqs/box_inverse.jpg
new file mode 100644
index 0000000000000000000000000000000000000000..dabcfb7b5d9d3ee97dd9736428bac4fb64f9d33c
GIT binary patch
literal 4570
zcmYLM1zc3k7Qeg8?$Wt{l8Z=p3zCw8z>?CXASETK!~)W(G)PN#xRih>AtfNaf^@0~
z0+Pa8zUTMe%F<3;+=h0N!66EgZdSxVQbY=f4!32mV(BB$WsG
zmk0h_sa}2o6ajn~3=V_h!{KlO0(?SZ1PL(_5%G0ODl!Bk9WxUn9RmaMCO-%A1}`fE
z1E&NRub_~qs3I`uG
z+cfKtUsMblxFOFZjRoN1?U}U_t5$U-eewYD$(wnCvU|^JEb{;*gAKLV{XFY0NboKC
z!}*Mo+_W--xt}h{Xfo-X8jN2@f|T`ePAnRHlOj!rfQ=v+N$5N6@Su2Dq+*(FYR#di
zsnljK@YN3{;Q`RL^msmnOH%eXNu&LC0dY_!?
z*`_b9?hEhQ@THg-*9N$>VaSseE
zQUvnlGr1@gqN7`sw~vg`-R{c!+YV}>s;`h{BU~;K*)biCJRjIrb;S@rOeS(1c1Bd@
z<74)V_A_`;4ma`E*+rbcEREUc13k~DZ~Mg>b2E|4)00VjgTP3Kgc0irqlpl`Jvzd{
zEvG%dgG2kR>kgiohKSyKGam#Zb{{duvjiA-of}`BL;%Z7^Q3(d3MLL
zQesEf?6?~eGfp0JIdcuVTF|WLTmt6%K^+&4FYCtzG`7Q9IDg78r1<0>eQNq)XPOn*
zZWxwv@MC!5RYv)YVYWzr3B#ym+hmvsbz*qZ`T10@*G6DEYRYJW+b}Gis&*st)3WiL
zhUBtH&h7oaK;I4NC(V(c){W-|C0AZ0ZSSyUcQ@on4ije$cVwcuF%jIugv9oRDbXyu
zxl=iFwR6nQ7vSPcpsycXgBd6Ode^93j7&@u`9=G%jZ>5VPiOB*^E*|ay$;t@TdL2l
z9~SwePvwE*^13UZ8M{MY#eixzF`QO4-I)9=qZ>O*pYSlJWltRXGsd?Jw4j7CX}N}v
zo6du9aUw~MixMyx4u!(txX{5h5CEp2WPu2g-w;;NvCgML%DV^SAq4Muu!^8w)GdoD
z>e{?4So#|}FgftvCpSv@<_Q<*#m{4yBERT@PEKA-gh
z2dZ}3EbB%!gQT#cThP?&o$G!3{95GehDzo
zd%ePMNMsNl%p}mA0z4H~xcVe+^~?0%s;(|0jZGrJV-U$@CO5xlmT>Kt2*^q-LxRk#
z__v>$mK&^geygPM=6IkJZKM0OP^#k9KC)$=!ATL0+P7WT6T~8gh{e#lS>#5IpgRqV2mJVpCu0b3*fzk{(uxmh8ettIQ5A3b)gG1>Ri
zD?&;eVPE%K;Hl%c8&&OARtp>4_Q%*BjXm}>J}^GelR=Cq=bl!c@}BZl$2b-Isc^l2
z>R@V4Gp(Gfe5O-)u;VSbHH7G$B!R6&i%At|MY(lmrp9*?_Q(6fnd!rK7aF_72pA^1
zY0ZuYYWCbs7NJnM^bA9D@i&8!c7A`rV>v^wT>4o?arkmSmBIeneA!5N>`6syMfUz~
zK0Zv8^geGDldT$jI#>e~A^VL);$XDjE~Oygt*`k_Gvf6BNn
zx;tUy>lxQRI%c}aoN}5x(d)(_syX7*C=e<=gUin3{GY!4E_+ym$5k?J>D&^eBWpB3
zga)M6&00Jzdg$+b7`y~13*v0sa<^SW`kfyvMo_;UaHlSEmEo_acBLtE2}}FhBBy)W
z;7WTMOB`d{D!dzK6DMve)A{@HFRC{`@FDFBRvnMA7bvgk2wT*t5lQ?|ISW=0wX4DH
zZI~1h-L2~eZkv$ZrJVPa%y`Y;o?C1E<0J~})Ke0WPnQNcmXcz^hCi`vIDKuIH(pBTu#5)vU%K9_Q{hr+9k{09B)dGo*zyIGZGKVLs
zmrCbvey7!4a|>Mb$SP+k&T>b1A{OQF6jW6`8w{1NRf>DHULd0}F9OD?rTi2p34=@0
zv*_Ig=L9{er(5kGxdEP(lk(E6FsFx2Te`(r4BThL4p7-y=i}pu8jG;5=Phl;H3e&f
zXAI+q`|iOla&Gi|7jZinP$N(_yeiW-||Q~~A&o*mP3X;|5v
zH{AMVx;eR#$v)uE|M|SCKqY<*2qcv1R|FT)g>>1QZYajA3G_x
zAOda)DY{WZykg_MZf`PXJj<9@F44fr@}M0{UNaXo*URVBaI
ztWJM$rTVvb16)`51I8f`9u9u4ign!m{SW?Fle@n_3I^xP*Nt_dmiGQF_2kGe+?dD@
zL|Z$%RYN_d4bA1pm^Ik4PmO^shAFw|cViYUDET$!o21Pt71bFIs#d>=>{k4htv-+)
z5hR6woIg?2ON^L9wI@+ONrH1Ebp^7V--B4b1kZZE>E;(x|wiSujZt)RLv$#!o@Quc8&+4LB@{6^&l;J&g
zOn2j}S{g@4UHG~L%p&mA?4~y;2+Pu(P;hK{Kb>12hzbl=eWr*9G?^FW-L)9*HpnY=
znEif;zz|aCd-{gd3TaYL4{OQ=KVeijtu||Mr>H{F9OEt6dFztgaLWw{_0R78btYj>
zT)wGbnYgCUFGkm>Z+hbgUh^9`8uDm~?H!SLAdD_+F})o5ly4+5RF0r11#+-1oj)DB
z`rFHC9c*mw0b{n6J%T2tI_DL}<#@#4g7{~kAzV*&1RP6+xF<;aDA*0o7Ix*BGpr&qQAuj+=bA{+
z@qtDhExJ&I9$Q*`jrDncPaf+pEgQ5af(dx1Zr58FrbNhYEkbQ=vvcI!a?YYi-HQl4
zR)`DuQw^2VqkN6eO``ZYHov}DPdsj*%EAg(=T=`MqXa$W=VJ=lX-DLY-1WG?0rZwd
ziC=AE7)^Pp{WENV6T`dR(uYn?aVov9_4=JiTp8tcYHjo_CCp8!74_tyq|rwhi?C$O
z>jy>WFst#&sY;LFmwQIh$9bazCmx^)lU^=CKtP?M)XctSARVOJtMJpi`%8WWqcMhz
z>^_yO7hBnr714u3usfLw
zvFRYojsSBZJlgtZeae8X75)#ipQMx4No_{zUWhV^Z
z_(m8wwKAaQXuI@P4O~fS2sey_gf8A+*rbuhSW|xd+^XC&Wvc30M#(-mC4k%?c)X83
z0-fDybyI@nHB*E=jhVXhW~K*!7Hv5vW3H-Fvq$`jvcQbeqA2Np=6b$ne#Iha5>F#x
z6sYxDw7bb@a{SxKQHN`yL^oUA>%T-OKsMS%U97d$xxO?!}-o9kR6|)X5;hM&eEunut{89G&nN;=IQl%7kgSi$?5X0kr~h7
z`<=Qhu38VkN;8~_kjR86y;hGdzd}`XG<~Y@AKnyaC~zmda+T)YT%&hE%z9A|L3__n
z2H4j!P8msBn=>QxJqxSFG)Caak;B`d%}a=8lryMZK86T&U8bDn_>4&DwG60RCW0f<
zgqXW5rkN0FuLJFySkXAU<9+S@#pnb~McaYp89Q@b>6k`-@2pjD3OICES#UgfA~aap
zq%+12&NOg6tOF`J^2CT8J1HwdbbLFj_Bnijme6BFR*mdQNVJ#RA}Vr87Rq&m1}POL
z=LD&5`}8yp0xOVv?JztDjAu|aDP-f8rvT&G2Dc@9D-qJ
Nly8{jxy0q{e*k;HJCFbX
literal 0
HcmV?d00001
diff --git a/doc/Eqs/box_inverse.tex b/doc/Eqs/box_inverse.tex
new file mode 100644
index 0000000000..68994f2868
--- /dev/null
+++ b/doc/Eqs/box_inverse.tex
@@ -0,0 +1,14 @@
+\documentclass[12pt]{article}
+
+\begin{document}
+
+\begin{eqnarray*}
+{\rm lx} &=& a \\
+{\rm xy} &=& b \cos{\gamma} \\
+{\rm xz} &=& c \cos{\beta}\\
+{\rm ly}^2 &=& b^2 - {\rm xy}^2 \\
+{\rm yz} &=& \frac{b*c \cos{\alpha} - {\rm xy}*{\rm xz}}{\rm ly} \\
+{\rm lz}^2 &=& c^2 - {\rm xz}^2 - {\rm yz}^2 \\
+\end{eqnarray*}
+
+\end{document}
diff --git a/doc/Section_howto.html b/doc/Section_howto.html
index adab09b06b..655adf939a 100644
--- a/doc/Section_howto.html
+++ b/doc/Section_howto.html
@@ -825,6 +825,10 @@ factors (xy,xz,yz) is as follows:
+The inverse relationship can be written as follows:
+
+
+
As discussed on the dump command doc page, when the BOX
BOUNDS for a snapshot is written to a dump file for a triclinic box,
an orthogonal bounding box which encloses the triclinic simulation box
@@ -1158,7 +1162,7 @@ discussed below, it can be referenced via the following bracket
notation, where ID in this case is the ID of a compute. The leading
"c_" would be replaced by "f_" for a fix, or "v_" for a variable:
-
+
| c_ID | entire scalar, vector, or array |
| c_ID[I] | one element of vector, one column of array |
| c_ID[I][J] | one element of array
@@ -1352,7 +1356,7 @@ data and scalar/vector/array data.
input, that could be an element of a vector or array. Likewise a
vector input could be a column of an array.
-
+
| Command | Input | Output | |
| thermo_style custom | global scalars | screen, log file | |
| dump custom | per-atom vectors | dump file | |
From dde38662d21d7d0c565f988cb9fd17b61e2f1dd2 Mon Sep 17 00:00:00 2001
From: sjplimp
Date: Mon, 23 May 2011 22:56:45 +0000
Subject: [PATCH 83/83] git-svn-id:
svn://svn.icms.temple.edu/lammps-ro/trunk@6215
f3b2605a-c512-4ea7-a41b-209d697bcdaa
---
src/GRANULAR/pair_gran_hooke_history.h | 2 +-
src/KSPACE/pppm.h | 2 +-
src/MOLECULE/atom_vec_angle.h | 2 +-
src/MOLECULE/atom_vec_full.h | 2 +-
src/MOLECULE/pair_lj_charmm_coul_charmm.h | 2 +-
src/atom_vec_atomic.h | 2 +-
src/atom_vec_charge.h | 2 +-
src/fix_nve.h | 1 +
8 files changed, 8 insertions(+), 7 deletions(-)
diff --git a/src/GRANULAR/pair_gran_hooke_history.h b/src/GRANULAR/pair_gran_hooke_history.h
index a27b6ce916..2bc2cf63cc 100644
--- a/src/GRANULAR/pair_gran_hooke_history.h
+++ b/src/GRANULAR/pair_gran_hooke_history.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairGranHookeHistory : public Pair {
public:
PairGranHookeHistory(class LAMMPS *);
- ~PairGranHookeHistory();
+ virtual ~PairGranHookeHistory();
virtual void compute(int, int);
virtual void settings(int, char **);
void coeff(int, char **);
diff --git a/src/KSPACE/pppm.h b/src/KSPACE/pppm.h
index 3577d6124c..f0ee75d12a 100644
--- a/src/KSPACE/pppm.h
+++ b/src/KSPACE/pppm.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PPPM : public KSpace {
public:
PPPM(class LAMMPS *, int, char **);
- ~PPPM();
+ virtual ~PPPM();
void init();
void setup();
void compute(int, int);
diff --git a/src/MOLECULE/atom_vec_angle.h b/src/MOLECULE/atom_vec_angle.h
index cfe83c716e..def4f67743 100644
--- a/src/MOLECULE/atom_vec_angle.h
+++ b/src/MOLECULE/atom_vec_angle.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class AtomVecAngle : public AtomVec {
public:
AtomVecAngle(class LAMMPS *, int, char **);
- ~AtomVecAngle() {}
+ virtual ~AtomVecAngle() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
diff --git a/src/MOLECULE/atom_vec_full.h b/src/MOLECULE/atom_vec_full.h
index a457c04660..8749c183de 100644
--- a/src/MOLECULE/atom_vec_full.h
+++ b/src/MOLECULE/atom_vec_full.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class AtomVecFull : public AtomVec {
public:
AtomVecFull(class LAMMPS *, int, char **);
- ~AtomVecFull() {}
+ virtual ~AtomVecFull() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
diff --git a/src/MOLECULE/pair_lj_charmm_coul_charmm.h b/src/MOLECULE/pair_lj_charmm_coul_charmm.h
index 970e03805c..2b3dae2274 100644
--- a/src/MOLECULE/pair_lj_charmm_coul_charmm.h
+++ b/src/MOLECULE/pair_lj_charmm_coul_charmm.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class PairLJCharmmCoulCharmm : public Pair {
public:
PairLJCharmmCoulCharmm(class LAMMPS *);
- ~PairLJCharmmCoulCharmm();
+ virtual ~PairLJCharmmCoulCharmm();
virtual void compute(int, int);
void settings(int, char **);
void coeff(int, char **);
diff --git a/src/atom_vec_atomic.h b/src/atom_vec_atomic.h
index 537065c553..34fb3779a7 100644
--- a/src/atom_vec_atomic.h
+++ b/src/atom_vec_atomic.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class AtomVecAtomic : public AtomVec {
public:
AtomVecAtomic(class LAMMPS *, int, char **);
- ~AtomVecAtomic() {}
+ virtual ~AtomVecAtomic() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
diff --git a/src/atom_vec_charge.h b/src/atom_vec_charge.h
index cfd6555acb..eb9fb91054 100644
--- a/src/atom_vec_charge.h
+++ b/src/atom_vec_charge.h
@@ -27,7 +27,7 @@ namespace LAMMPS_NS {
class AtomVecCharge : public AtomVec {
public:
AtomVecCharge(class LAMMPS *, int, char **);
- ~AtomVecCharge() {}
+ virtual ~AtomVecCharge() {}
void grow(int);
void grow_reset();
void copy(int, int, int);
diff --git a/src/fix_nve.h b/src/fix_nve.h
index 8ada964608..2be7a671e8 100644
--- a/src/fix_nve.h
+++ b/src/fix_nve.h
@@ -27,6 +27,7 @@ namespace LAMMPS_NS {
class FixNVE : public Fix {
public:
FixNVE(class LAMMPS *, int, char **);
+ virtual ~FixNVE() {}
int setmask();
virtual void init();
virtual void initial_integrate(int);
|