From 4c2439be9e6054ac9ddecaddda7526f18aec70be Mon Sep 17 00:00:00 2001 From: Lu Date: Wed, 1 Apr 2020 15:55:14 +0800 Subject: [PATCH 1/4] fix bug of load transform --- source/train/CMakeLists.txt | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/train/CMakeLists.txt b/source/train/CMakeLists.txt index 8be6b6c819..84a0848a77 100644 --- a/source/train/CMakeLists.txt +++ b/source/train/CMakeLists.txt @@ -2,7 +2,7 @@ configure_file("RunOptions.py.in" "${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py" @ONLY) -file(GLOB LIB_PY main.py common.py env.py compat.py calculator.py Network.py Deep*.py Data.py DataSystem.py Model*.py Descrpt*.py Fitting.py Loss.py LearningRate.py Trainer.py TabInter.py EwaldRecp.py DataModifier.py ${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py) +file(GLOB LIB_PY main.py common.py env.py compat.py calculator.py Network.py Deep*.py Data.py DataModifier.py DataSystem.py Model*.py Descrpt*.py Fitting.py Loss.py LearningRate.py Trainer.py TabInter.py EwaldRecp.py DataModifier.py ${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py transform.py print_old_model.py) file(GLOB CLS_PY Local.py Slurm.py) From 0e42066b8852856c02aea6af8f570cc1ab577d1d Mon Sep 17 00:00:00 2001 From: Lu Date: Thu, 2 Apr 2020 11:01:18 +0800 Subject: [PATCH 2/4] Increase the max neighbor size of an atom from 256 to 1024 --- source/lib/src/NNPInter.cc | 2 +- source/op/cuda/descrpt_se_a.cu | 4 ++-- source/op/descrpt_se_a_gpu.cc | 2 -- 3 files changed, 3 insertions(+), 5 deletions(-) diff --git a/source/lib/src/NNPInter.cc b/source/lib/src/NNPInter.cc index 58042019c7..f4f39945ff 100644 --- a/source/lib/src/NNPInter.cc +++ b/source/lib/src/NNPInter.cc @@ -3,7 +3,7 @@ #include "SimulationRegion.h" #include -#define MAGIC_NUMBER 256 +#define MAGIC_NUMBER 1024 #ifdef USE_CUDA_TOOLKIT #include "cuda_runtime.h" diff --git a/source/op/cuda/descrpt_se_a.cu b/source/op/cuda/descrpt_se_a.cu index 39434306f2..4b309522fa 100644 --- a/source/op/cuda/descrpt_se_a.cu +++ b/source/op/cuda/descrpt_se_a.cu @@ -18,7 +18,7 @@ limitations under the License. #include #include -#define MAGIC_NUMBER 256 +#define MAGIC_NUMBER 1024 #ifdef HIGH_PREC typedef double VALUETYPE; @@ -326,7 +326,7 @@ void DescrptSeALauncher(const VALUETYPE* coord, i_idx ); const int ITEMS_PER_THREAD = 4; - const int BLOCK_THREADS = 64; + const int BLOCK_THREADS = MAGIC_NUMBER / ITEMS_PER_THREAD; // BlockSortKernel<<>> ( BlockSortKernel <<>> (key, key + nloc * MAGIC_NUMBER); diff --git a/source/op/descrpt_se_a_gpu.cc b/source/op/descrpt_se_a_gpu.cc index 93c83016fb..a4b6b3512a 100644 --- a/source/op/descrpt_se_a_gpu.cc +++ b/source/op/descrpt_se_a_gpu.cc @@ -7,7 +7,6 @@ #include "tensorflow/core/framework/shape_inference.h" using namespace tensorflow; // NOLINT(build/namespaces) -#define MAGIC_NUMBER 256 #ifdef HIGH_PREC typedef double VALUETYPE ; @@ -201,7 +200,6 @@ class DescrptSeAOp : public OpKernel { cudaErrcheck(cudaMemcpy(&(array_longlong), 20 + mesh_tensor.flat().data(), sizeof(unsigned long long *), cudaMemcpyDeviceToHost)); cudaErrcheck(cudaMemcpy(&(array_double), 24 + mesh_tensor.flat().data(), sizeof(compute_t *), cudaMemcpyDeviceToHost)); - // cudaErrcheck(cudaMemcpy(jlist, host_jlist, sizeof(int) * nloc * MAGIC_NUMBER, cudaMemcpyHostToDevice)); // Launch computation for (int II = 0; II < nsamples; II++) { DescrptSeALauncher(coord_tensor.matrix().data() + II * (nall * 3), // related to the kk argument From 62cc79a0bb9e6583b7d686367adf39b357f75e4b Mon Sep 17 00:00:00 2001 From: Lu Date: Thu, 2 Apr 2020 16:54:12 +0800 Subject: [PATCH 3/4] Add boundary check if nnei is larger than 1024 --- source/op/descrpt_se_a_gpu.cc | 3 ++- source/train/CMakeLists.txt | 2 +- 2 files changed, 3 insertions(+), 2 deletions(-) diff --git a/source/op/descrpt_se_a_gpu.cc b/source/op/descrpt_se_a_gpu.cc index a4b6b3512a..09e23d010d 100644 --- a/source/op/descrpt_se_a_gpu.cc +++ b/source/op/descrpt_se_a_gpu.cc @@ -158,7 +158,8 @@ class DescrptSeAOp : public OpKernel { OP_REQUIRES (context, (ntypes == int(sel_a.size())), errors::InvalidArgument ("number of types should match the length of sel array")); OP_REQUIRES (context, (ntypes == int(sel_r.size())), errors::InvalidArgument ("number of types should match the length of sel array")); - + OP_REQUIRES (context, (nnei > 1024), errors::InvalidArgument ("Assert failed, max neighbor size of atom(nnei) " + std::to_string(nnei) + " is larger than 1024!, which currently is not supported by deepmd-kit.")); + // Create output tensors TensorShape descrpt_shape ; descrpt_shape.AddDim (nsamples); diff --git a/source/train/CMakeLists.txt b/source/train/CMakeLists.txt index 84a0848a77..1875d2097c 100644 --- a/source/train/CMakeLists.txt +++ b/source/train/CMakeLists.txt @@ -2,7 +2,7 @@ configure_file("RunOptions.py.in" "${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py" @ONLY) -file(GLOB LIB_PY main.py common.py env.py compat.py calculator.py Network.py Deep*.py Data.py DataModifier.py DataSystem.py Model*.py Descrpt*.py Fitting.py Loss.py LearningRate.py Trainer.py TabInter.py EwaldRecp.py DataModifier.py ${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py transform.py print_old_model.py) +file(GLOB LIB_PY main.py common.py env.py compat.py calculator.py Network.py Deep*.py Data.py DataSystem.py Model*.py Descrpt*.py Fitting.py Loss.py LearningRate.py Trainer.py TabInter.py EwaldRecp.py DataModifier.py ${CMAKE_CURRENT_BINARY_DIR}/RunOptions.py transform.py) file(GLOB CLS_PY Local.py Slurm.py) From 3a38dc4737c264fa513ed4c8d6572848da47a8c7 Mon Sep 17 00:00:00 2001 From: Lu Date: Thu, 2 Apr 2020 16:57:18 +0800 Subject: [PATCH 4/4] Add boundary check if nnei is larger than 1024 --- source/op/descrpt_se_a_gpu.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/source/op/descrpt_se_a_gpu.cc b/source/op/descrpt_se_a_gpu.cc index 09e23d010d..70dd9c7751 100644 --- a/source/op/descrpt_se_a_gpu.cc +++ b/source/op/descrpt_se_a_gpu.cc @@ -158,7 +158,7 @@ class DescrptSeAOp : public OpKernel { OP_REQUIRES (context, (ntypes == int(sel_a.size())), errors::InvalidArgument ("number of types should match the length of sel array")); OP_REQUIRES (context, (ntypes == int(sel_r.size())), errors::InvalidArgument ("number of types should match the length of sel array")); - OP_REQUIRES (context, (nnei > 1024), errors::InvalidArgument ("Assert failed, max neighbor size of atom(nnei) " + std::to_string(nnei) + " is larger than 1024!, which currently is not supported by deepmd-kit.")); + OP_REQUIRES (context, (nnei <= 1024), errors::InvalidArgument ("Assert failed, max neighbor size of atom(nnei) " + std::to_string(nnei) + " is larger than 1024!, which currently is not supported by deepmd-kit.")); // Create output tensors TensorShape descrpt_shape ;