merge recent changes on devel back to master

.travis.yml

@@ -83,7 +83,7 @@ jobs:
         - CXX=g++-7
         - TENSORFLOW_VERSION=2.1
       install:
-        - python -m pip install twine cibuildwheel==1.1.0 scikit-build
+        - python -m pip install twine cibuildwheel==1.1.0 scikit-build setuptools_scm
       script:
         - python -m cibuildwheel --output-dir wheelhouse
         - python setup.py sdist

pyproject.toml

@@ -1,3 +1,3 @@
 [build-system]
-requires = ["setuptools", "wheel", "scikit-build", "cmake", "ninja", "m2r"]
+requires = ["setuptools", "setuptools_scm", "wheel", "scikit-build", "cmake", "ninja", "m2r"]
 

setup.py

@@ -1,9 +1,21 @@
 from skbuild import setup
 from skbuild.exceptions import SKBuildError
 from skbuild.cmaker import get_cmake_version
+from setuptools_scm import get_version
 from packaging.version import LegacyVersion
 from os import path, makedirs
-import imp
+import imp, sys, platform
+
+def get_dp_install_path() :
+    site_packages_path = path.join(path.dirname(path.__file__), 'site-packages')
+    dp_scm_version     = get_version(root="./", relative_to=__file__)
+    python_version     = 'py' + str(sys.version_info.major + sys.version_info.minor * 0.1)
+    os_info            = sys.platform
+    machine_info       = platform.machine()
+    dp_pip_install_path    = site_packages_path + '/deepmd'
+    dp_setup_install_path    = site_packages_path + '/deepmd_kit-' + dp_scm_version + '-' + python_version + '-' + os_info + '-' + machine_info + '.egg/deepmd'
+
+    return dp_pip_install_path, dp_setup_install_path
 
 readme_file = path.join(path.dirname(path.abspath(__file__)), 'README.md')
 try:
@@ -34,6 +46,8 @@
 except OSError:
     pass
 
+dp_pip_install_path, dp_setup_install_path = get_dp_install_path()
+
 setup(
     name="deepmd-kit",
     setup_requires=setup_requires,
@@ -56,6 +70,8 @@
                 '-DBUILD_PY_IF:BOOL=TRUE', 
                 '-DBUILD_CPP_IF:BOOL=FALSE',
                 '-DFLOAT_PREC:STRING=high',
+                '-DDP_PIP_INSTALL_PATH=%s' % dp_pip_install_path,
+                '-DDP_SETUP_INSTALL_PATH=%s' % dp_setup_install_path,
     ],
     cmake_source_dir='source',
     cmake_minimum_required_version='3.0',

source/CMakeLists.txt

@@ -186,7 +186,7 @@ if (BUILD_CPP_IF)
   set (LIB_DEEPMD_OP		"deepmd_op")
   if (USE_CUDA_TOOLKIT)
     set (LIB_DEEPMD_OP_CUDA		"deepmd_op_cuda")
-  else()
+  else ()
     set (LIB_DEEPMD_OP_CUDA		"deepmd_op")
   endif()
   if (CMAKE_CXX_COMPILER_VERSION VERSION_GREATER_EQUAL 4.9)

source/lib/include/NNPInter.h

@@ -93,10 +93,8 @@ class NNPInter
   compute_t *array_double;
   InternalNeighborList nlist;
   NNPAtomMap<VALUETYPE> nnpmap;
-  unsigned long long *array_longlong;
-  int *ilist, *jrange, *jlist, *array_int;
+  int *ilist, *jrange, *jlist;
   int ilist_size, jrange_size, jlist_size;
-  int arr_int_size, arr_ll_size, arr_dou_size;
 
   // function used for neighbor list copy
   vector<int> get_sel_a() const;
@@ -191,13 +189,10 @@ class NNPInterModelDevi
   vector<vector<int> > sec;
   InternalNeighborList nlist;
   NNPAtomMap<VALUETYPE> nnpmap;
-  unsigned long long *array_longlong;
-  int max_sec_size = 0, max_sec_back = 0;
-  int *ilist, *jrange, *jlist, *array_int;
-  int ilist_size, jrange_size, jlist_size, arr_int_size, arr_ll_size, arr_dou_size;
+  int *ilist, *jrange, *jlist;
+  int ilist_size, jrange_size, jlist_size;
 
   // function used for nborlist copy
-  void get_max_sec();
   vector<vector<int> > get_sel() const;
   void cum_sum(const std::vector<std::vector<int32> > n_sel);
 #ifdef USE_CUDA_TOOLKIT

source/lib/include/SimulationRegion_Impl.h

@@ -5,6 +5,7 @@
 #include <iostream>
 #include <limits>
 #include <typeinfo>
+#include <stdexcept>
 
 using namespace std;
 
@@ -500,6 +501,9 @@ computeVolume()
       boxt[0*3+1] * (boxt[1*3+0]*boxt[2*3+2] - boxt[2*3+0]*boxt[1*3+2]) +
       boxt[0*3+2] * (boxt[1*3+0]*boxt[2*3+1] - boxt[2*3+0]*boxt[1*3+1]);
   volumei = static_cast<double>(1.)/volume;
+  if (volume < 0) {
+    throw std::runtime_error("Negative volume detected. Please make sure the simulation cell obeys the right-hand rule.");
+  }
 }
 
 template<typename VALUETYPE>

source/lib/include/common.h

@@ -144,9 +144,6 @@ session_input_tensors (vector<std::pair<string, Tensor>>& input_tensors,
 		       const int                        * ilist, 
 		       const int                        * jrange,
 		       const int                        * jlist,
-		       int                              * array_int,
-		       unsigned long long               * array_longlong, 
-		       double                           * array_double,
 		       const vector<VALUETYPE>		& fparam_,
 		       const vector<VALUETYPE>	        & aparam_,
 		       const NNPAtomMap<VALUETYPE>      & nnpmap,

source/lib/src/NNPInter.cc

@@ -3,7 +3,6 @@
 #include "SimulationRegion.h"
 #include <stdexcept>	
 
-#define MAGIC_NUMBER 1024
 
 #ifdef  USE_CUDA_TOOLKIT
 #include "cuda_runtime.h"
@@ -14,7 +13,7 @@
 #define cudaErrcheck(res) { cudaAssert((res), __FILE__, __LINE__); }
 inline void cudaAssert(cudaError_t code, const char *file, int line, bool abort=true)
 {
-    if (code != cudaSuccess) 
+    if (code != cudaSuccess)
     {
         fprintf(stderr,"cuda assert: %s %s %d\n", cudaGetErrorString(code), file, line);
         if (abort) exit(code);
@@ -273,33 +272,19 @@ NNPInter::~NNPInter() {
         cudaErrcheck(cudaFree(ilist));
         cudaErrcheck(cudaFree(jrange));
         cudaErrcheck(cudaFree(jlist));
-        cudaErrcheck(cudaFree(array_int));
-        cudaErrcheck(cudaFree(array_longlong));
-        cudaErrcheck(cudaFree(array_double));
     }
     #endif
 }
 
 #ifdef USE_CUDA_TOOLKIT
 void NNPInter::update_nbor(const InternalNeighborList & nlist, const int nloc) {
     if (!init_nbor) {
-        sec_a = cum_sum(get_sel_a());
         cudaErrcheck(cudaMalloc((void**)&ilist, sizeof(int) * nlist.ilist.size()));
         cudaErrcheck(cudaMalloc((void**)&jrange, sizeof(int) * nlist.jrange.size()));
         cudaErrcheck(cudaMalloc((void**)&jlist, sizeof(int) * nlist.jlist.size()));
-        cudaErrcheck(cudaMalloc((void**)&array_int, sizeof(int) * (sec_a.size() + nloc * sec_a.size() + nloc)));
-        cudaErrcheck(cudaMalloc((void**)&array_longlong, sizeof(unsigned long long) * nloc * MAGIC_NUMBER * 2));
-        #ifdef HIGH_PREC
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * sec_a.back() * 3));
-        #else
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * sec_a.back() * 3));
-        #endif
         ilist_size = nlist.ilist.size();
         jrange_size = nlist.jrange.size();
         jlist_size = nlist.jlist.size();
-        arr_int_size = sec_a.size() + nloc * sec_a.size() + nloc;
-        arr_ll_size = nloc * MAGIC_NUMBER * 2;
-        arr_dou_size = nloc * sec_a.back() * 3;
         init_nbor = true;
     }
     if (ilist_size < nlist.ilist.size()) {
@@ -317,25 +302,7 @@ void NNPInter::update_nbor(const InternalNeighborList & nlist, const int nloc) {
         cudaErrcheck(cudaMalloc((void**)&jlist, sizeof(int) * nlist.jlist.size()));
         jlist_size = nlist.jlist.size();
     }
-    if (arr_int_size < sec_a.size() + nloc * sec_a.size() + nloc) {
-        cudaErrcheck(cudaFree(array_int));
-        cudaErrcheck(cudaMalloc((void**)&array_int, sizeof(int) * (sec_a.size() + nloc * sec_a.size() + nloc)));
-        arr_int_size = sec_a.size() + nloc * sec_a.size() + nloc;
-    }
-    if (arr_ll_size < nloc * MAGIC_NUMBER * 2) {
-        cudaErrcheck(cudaFree(array_longlong));
-        cudaErrcheck(cudaMalloc((void**)&array_longlong, sizeof(unsigned long long) * nloc * MAGIC_NUMBER * 2));
-        arr_ll_size = nloc * MAGIC_NUMBER * 2;
-    }
-    if (arr_dou_size < nloc * sec_a.back() * 3) {
-        cudaErrcheck(cudaFree(array_double));
-        #ifdef HIGH_PREC
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * sec_a.back() * 3));
-        #else
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * sec_a.back() * 3));
-        #endif
-        arr_dou_size = nloc * sec_a.back() * 3;
-    }
+
     cudaErrcheck(cudaMemcpy(ilist, &nlist.ilist[0], sizeof(int) * nlist.ilist.size(), cudaMemcpyHostToDevice));
     cudaErrcheck(cudaMemcpy(jrange, &nlist.jrange[0], sizeof(int) * nlist.jrange.size(), cudaMemcpyHostToDevice));
     cudaErrcheck(cudaMemcpy(jlist, &nlist.jlist[0], sizeof(int) * nlist.jlist.size(), cudaMemcpyHostToDevice));
@@ -378,14 +345,10 @@ init (const string & model, const int & gpu_rank)
   if (dfparam < 0) dfparam = 0;
   if (daparam < 0) daparam = 0;
   inited = true;
-
+  
   init_nbor = false;
-  array_int = NULL;
-  array_double = NULL;
-  array_longlong = NULL;
   ilist = NULL; jrange = NULL; jlist = NULL;
   ilist_size = 0; jrange_size = 0; jlist_size = 0;
-  arr_int_size = 0; arr_ll_size = 0; arr_dou_size = 0;
 }
 #else
 void
@@ -415,12 +378,8 @@ init (const string & model, const int & gpu_rank)
   inited = true;
 
   init_nbor = false;
-  array_int = NULL;
-  array_double = NULL;
-  array_longlong = NULL;
   ilist = NULL; jrange = NULL; jlist = NULL;
   ilist_size = 0; jrange_size = 0; jlist_size = 0;
-  arr_int_size = 0; arr_ll_size = 0; arr_dou_size = 0;
 }
 #endif
 
@@ -602,7 +561,7 @@ compute_inner (ENERGYTYPE &			dener,
     }
 
     #ifdef USE_CUDA_TOOLKIT
-        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, array_int, array_longlong, array_double, fparam, aparam, nnpmap, nghost);
+        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, fparam, aparam, nnpmap, nghost);
     #else
         int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, nlist, fparam, aparam, nnpmap, nghost);
     #endif
@@ -669,7 +628,7 @@ compute (ENERGYTYPE &			dener,
     }
 
     #ifdef USE_CUDA_TOOLKIT
-        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, array_int, array_longlong, array_double, fparam, aparam, nnpmap, nghost);
+        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, fparam, aparam, nnpmap, nghost);
     #else
         int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, nlist, fparam, aparam, nnpmap, nghost);
     #endif
@@ -710,9 +669,6 @@ NNPInterModelDevi::~NNPInterModelDevi() {
         cudaErrcheck(cudaFree(ilist));
         cudaErrcheck(cudaFree(jrange));
         cudaErrcheck(cudaFree(jlist));
-        cudaErrcheck(cudaFree(array_int));
-        cudaErrcheck(cudaFree(array_longlong));
-        cudaErrcheck(cudaFree(array_double));
     }
 #endif
 }
@@ -761,14 +717,10 @@ init (const vector<string> & models, const int & gpu_rank)
   // cell_size = rcut;
   // ntypes = get_ntypes();
   inited = true;
-
+  
   init_nbor = false;
-  array_int = NULL;
-  array_double = NULL;
-  array_longlong = NULL;
   ilist = NULL; jrange = NULL; jlist = NULL;
   ilist_size = 0; jrange_size = 0; jlist_size = 0;
-  arr_int_size = 0; arr_ll_size = 0; arr_dou_size = 0;
 }
 #else
 void
@@ -798,14 +750,10 @@ init (const vector<string> & models, const int & gpu_rank)
   // cell_size = rcut;
   // ntypes = get_ntypes();
   inited = true;
-
+  
   init_nbor = false;
-  array_int = NULL;
-  array_double = NULL;
-  array_longlong = NULL;
   ilist = NULL; jrange = NULL; jlist = NULL;
   ilist_size = 0; jrange_size = 0; jlist_size = 0;
-  arr_int_size = 0; arr_ll_size = 0; arr_dou_size = 0;
 }
 #endif
 
@@ -873,40 +821,18 @@ cum_sum (const std::vector<std::vector<int32> > n_sel)
     }
 }
 
-void  
-NNPInterModelDevi::
-get_max_sec() 
-{
-    for (int ii = 0; ii < numb_models; ii++) {
-        this->max_sec_size = max_sec_size < sec[ii].size() ? sec[ii].size() : max_sec_size;
-        this->max_sec_back = max_sec_back < sec[ii].back() ? sec[ii].back() : max_sec_back;
-    }
-}
-
 #ifdef USE_CUDA_TOOLKIT
 void
 NNPInterModelDevi::
 update_nbor(const InternalNeighborList & nlist, const int nloc) 
 {
     if (!init_nbor) {
-        cum_sum(get_sel());
-        get_max_sec();
         cudaErrcheck(cudaMalloc((void**)&ilist, sizeof(int) * nlist.ilist.size()));
         cudaErrcheck(cudaMalloc((void**)&jrange, sizeof(int) * nlist.jrange.size()));
         cudaErrcheck(cudaMalloc((void**)&jlist, sizeof(int) * nlist.jlist.size()));
-        cudaErrcheck(cudaMalloc((void**)&array_int, sizeof(int) * (max_sec_size + nloc * max_sec_size + nloc)));
-        cudaErrcheck(cudaMalloc((void**)&array_longlong, sizeof(unsigned long long) * nloc * MAGIC_NUMBER * 2));
-        #ifdef HIGH_PREC
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * max_sec_back * 3));
-        #else
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * max_sec_back * 3));
-        #endif
         ilist_size = nlist.ilist.size();
         jrange_size = nlist.jrange.size();
         jlist_size = nlist.jlist.size();
-        arr_int_size = max_sec_size + nloc * max_sec_size + nloc;
-        arr_ll_size = nloc * MAGIC_NUMBER * 2;
-        arr_dou_size = nloc * max_sec_back * 3;
         init_nbor = true;
     }
     if (ilist_size < nlist.ilist.size()) {
@@ -924,25 +850,7 @@ update_nbor(const InternalNeighborList & nlist, const int nloc)
         cudaErrcheck(cudaMalloc((void**)&jlist, sizeof(int) * nlist.jlist.size()));
         jlist_size = nlist.jlist.size();
     }
-    if (arr_int_size < max_sec_size + nloc * max_sec_size + nloc) {
-        cudaErrcheck(cudaFree(array_int));
-        cudaErrcheck(cudaMalloc((void**)&array_int, sizeof(int) * (max_sec_size + nloc * max_sec_size + nloc)));
-        arr_int_size = max_sec_size + nloc * max_sec_size + nloc;
-    }
-    if (arr_ll_size < nloc * MAGIC_NUMBER * 2) {
-        cudaErrcheck(cudaFree(array_longlong));
-        cudaErrcheck(cudaMalloc((void**)&array_longlong, sizeof(unsigned long long) * nloc * MAGIC_NUMBER * 2));
-        arr_ll_size = nloc * MAGIC_NUMBER * 2;
-    }
-    if (arr_dou_size < nloc * max_sec_back * 3) {
-        cudaErrcheck(cudaFree(array_double));
-        #ifdef HIGH_PREC
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * max_sec_back * 3));
-        #else
-            cudaErrcheck(cudaMalloc((void**)&array_double, sizeof(compute_t) * nloc * max_sec_back * 3));
-        #endif
-        arr_dou_size = nloc * max_sec_back * 3;
-    }
+
     cudaErrcheck(cudaMemcpy(ilist, &nlist.ilist[0], sizeof(int) * nlist.ilist.size(), cudaMemcpyHostToDevice));
     cudaErrcheck(cudaMemcpy(jrange, &nlist.jrange[0], sizeof(int) * nlist.jrange.size(), cudaMemcpyHostToDevice));
     cudaErrcheck(cudaMemcpy(jlist, &nlist.jlist[0], sizeof(int) * nlist.jlist.size(), cudaMemcpyHostToDevice));
@@ -1044,7 +952,7 @@ compute (vector<ENERGYTYPE> &		all_energy,
 
     }
     #ifdef USE_CUDA_TOOLKIT
-        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, array_int, array_longlong, array_double, fparam, aparam, nnpmap, nghost);
+        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, fparam, aparam, nnpmap, nghost);
     #else
         int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, nlist, fparam, aparam, nnpmap, nghost);
     #endif
@@ -1094,7 +1002,7 @@ compute (vector<ENERGYTYPE> &			all_energy,
 
     }
     #ifdef USE_CUDA_TOOLKIT
-        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, array_int, array_longlong, array_double, fparam, aparam, nnpmap, nghost);
+        int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, ilist, jrange, jlist, fparam, aparam, nnpmap, nghost);
     #else
         int ret = session_input_tensors (input_tensors, dcoord_, ntypes, datype_, dbox, nlist, fparam, aparam, nnpmap, nghost);
     #endif