Skip to content

increase max neighbor size from 256 to 1024 #212

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
amcadmus merged 6 commits into from
Apr 2, 2020
Merged

increase max neighbor size from 256 to 1024 #212

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
86c7fcb
Merge pull request #13 from deepmodeling/devel
denghuilu Mar 30, 2020
28dc891
Merge pull request #14 from deepmodeling/devel
denghuilu Apr 1, 2020
4c2439b
fix bug of load transform
Apr 1, 2020
0e42066
Increase the max neighbor size of an atom from 256 to 1024
Apr 2, 2020
62cc79a
Add boundary check if nnei is larger than 1024
Apr 2, 2020
3a38dc4
Add boundary check if nnei is larger than 1024
Apr 2, 2020
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
2 changes: 1 addition & 1 deletion source/lib/src/NNPInter.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3,7 +3,7 @@
#include "SimulationRegion.h"
#include <stdexcept>

#define MAGIC_NUMBER 256
#define MAGIC_NUMBER 1024

#ifdef USE_CUDA_TOOLKIT
#include "cuda_runtime.h"
Expand Down
4 changes: 2 additions & 2 deletions source/op/cuda/descrpt_se_a.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -18,7 +18,7 @@ limitations under the License.
#include <cub/block/block_radix_sort.cuh>
#include <cuda_runtime.h>

#define MAGIC_NUMBER 256
#define MAGIC_NUMBER 1024

#ifdef HIGH_PREC
typedef double VALUETYPE;
Expand Down Expand Up @@ -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<NeighborInfo, BLOCK_THREADS, ITEMS_PER_THREAD><<<g_grid_size, BLOCK_THREADS>>> (
BlockSortKernel<int_64, BLOCK_THREADS, ITEMS_PER_THREAD> <<<nloc, BLOCK_THREADS>>> (key, key + nloc * MAGIC_NUMBER);

Expand Down
5 changes: 2 additions & 3 deletions source/op/descrpt_se_a_gpu.cc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 ;
Expand Down Expand Up @@ -159,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);
Expand Down Expand Up @@ -201,7 +201,6 @@ class DescrptSeAOp : public OpKernel {
cudaErrcheck(cudaMemcpy(&(array_longlong), 20 + mesh_tensor.flat<int>().data(), sizeof(unsigned long long *), cudaMemcpyDeviceToHost));
cudaErrcheck(cudaMemcpy(&(array_double), 24 + mesh_tensor.flat<int>().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<VALUETYPE>().data() + II * (nall * 3), // related to the kk argument
Expand Down
2 changes: 1 addition & 1 deletion source/train/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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 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)

Expand Down