diff --git a/.github/workflows/mirror_gitee.yml b/.github/workflows/mirror_gitee.yml index b10e2be62e..0d7205a03e 100644 --- a/.github/workflows/mirror_gitee.yml +++ b/.github/workflows/mirror_gitee.yml @@ -1,13 +1,23 @@ name: Mirror to Gitee Repo +# Ref creation is also delivered as a push event, but deletion may only emit a +# delete event. Keep both triggers so removed branches and tags are pruned from +# the mirror without launching duplicate mirrors for ref creation. on: push: - branches-ignore: - - "copilot/**" - - "dependabot/**" - - "pre-commit-ci-update-config" + # Use positive patterns with exclusions so tag pushes can be listed below. + # `branches-ignore` by itself disables all tag-triggered runs. + branches: + - "**" + # GitHub creates these ephemeral refs while testing a merge queue entry. + # They are neither source branches nor refs that Gitee should receive. + - "!gh-readonly-queue/**" + - "!copilot/**" + - "!dependabot/**" + - "!pre-commit-ci-update-config" + tags: + - "**" delete: - create: # Ensures that only one mirror task will run at a time. concurrency: @@ -15,8 +25,14 @@ concurrency: jobs: git-mirror: - if: github.repository_owner == 'deepmodeling' - runs-on: ubuntu-slim + # Delete events use the default branch as `github.ref`, so inspect the + # payload ref to suppress only ephemeral merge-queue branch deletions. + if: >- + github.repository_owner == 'deepmodeling' && + (github.event_name != 'delete' || + !startsWith(github.event.ref, 'gh-readonly-queue/')) + # This is a Docker action; ubuntu-slim does not provide a Docker daemon. + runs-on: ubuntu-latest steps: - uses: wearerequired/git-mirror-action@v1 env: diff --git a/source/lib/include/neighbor_stat.h b/source/lib/include/neighbor_stat.h deleted file mode 100644 index 79e241f783..0000000000 --- a/source/lib/include/neighbor_stat.h +++ /dev/null @@ -1,18 +0,0 @@ -// SPDX-License-Identifier: LGPL-3.0-or-later -#include "neighbor_list.h" - -#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM - -namespace deepmd { -template -void neighbor_stat_gpu(const FPTYPE* coord, - const int* type, - const int nloc, - const deepmd::InputNlist& gpu_nlist, - int* max_nbor_size, - FPTYPE* min_nbor_dist, - const int ntypes, - const int MAX_NNEI); -} // namespace deepmd - -#endif diff --git a/source/lib/src/gpu/neighbor_stat.cu b/source/lib/src/gpu/neighbor_stat.cu deleted file mode 100644 index ef7d3c5f8a..0000000000 --- a/source/lib/src/gpu/neighbor_stat.cu +++ /dev/null @@ -1,103 +0,0 @@ -#include - -#include "device.h" -#include "neighbor_list.h" - -template -__global__ void neighbor_stat_g(const FPTYPE* coord, - const int* type, - const int nloc, - const int* ilist, - int** firstneigh, - const int* numneigh, - int* max_nbor_size, - FPTYPE* min_nbor_dist, - const int ntypes, - const int MAX_NNEI) { - int ithread = blockIdx.x * blockDim.x + threadIdx.x; - int ii = ithread / MAX_NNEI; - int jj = ithread % MAX_NNEI; - // assume the same block has the same ii - __shared__ int cache[TPB]; - cache[threadIdx.x] = 0; - if (ii >= nloc) { - return; - } - int idx_i = ilist[ii]; - if (type[idx_i] < 0) { - // set all to 10000 - min_nbor_dist[ii * MAX_NNEI + jj] = INFINITY; - return; // virtual atom - } - if (jj < numneigh[ii]) { - int idx_j = firstneigh[ii][jj]; - int type_j = type[idx_j]; - if (type_j < 0) { - min_nbor_dist[ii * MAX_NNEI + jj] = INFINITY; - return; // virtual atom - } - __syncthreads(); - FPTYPE rij[3] = {coord[idx_j * 3 + 0] - coord[idx_i * 3 + 0], - coord[idx_j * 3 + 1] - coord[idx_i * 3 + 1], - coord[idx_j * 3 + 2] - coord[idx_i * 3 + 2]}; - // we do not need to use the real index - // we do not need to do slow sqrt for every dist; instead do sqrt in the - // final - min_nbor_dist[ii * MAX_NNEI + jj] = - rij[0] * rij[0] + rij[1] * rij[1] + rij[2] * rij[2]; - - // atomicAdd(max_nbor_size + ii * ntypes + type_j, 1); - // See https://www.cnblogs.com/neopenx/p/4705320.html - atomicAdd(&cache[type_j], 1); - __syncthreads(); - if (threadIdx.x < ntypes) { - atomicAdd(&max_nbor_size[ii * ntypes + threadIdx.x], cache[threadIdx.x]); - } - } else { - // set others to 10000 - min_nbor_dist[ii * MAX_NNEI + jj] = INFINITY; - } -} - -namespace deepmd { - -template -void neighbor_stat_gpu(const FPTYPE* coord, - const int* type, - const int nloc, - const deepmd::InputNlist& gpu_nlist, - int* max_nbor_size, - FPTYPE* min_nbor_dist, - const int ntypes, - const int MAX_NNEI) { - DPErrcheck(gpuGetLastError()); - DPErrcheck(gpuDeviceSynchronize()); - - DPErrcheck(gpuMemset(max_nbor_size, 0, sizeof(int) * int_64(nloc) * ntypes)); - const int nblock_loc = (nloc * MAX_NNEI + TPB - 1) / TPB; - neighbor_stat_g<<>>( - coord, type, nloc, gpu_nlist.ilist, gpu_nlist.firstneigh, - gpu_nlist.numneigh, max_nbor_size, min_nbor_dist, ntypes, MAX_NNEI); - - DPErrcheck(gpuGetLastError()); - DPErrcheck(gpuDeviceSynchronize()); -} - -template void neighbor_stat_gpu(const float* coord, - const int* type, - const int nloc, - const deepmd::InputNlist& gpu_nlist, - int* max_nbor_size, - float* min_nbor_dist, - const int ntypes, - const int MAX_NNEI); - -template void neighbor_stat_gpu(const double* coord, - const int* type, - const int nloc, - const deepmd::InputNlist& gpu_nlist, - int* max_nbor_size, - double* min_nbor_dist, - const int ntypes, - const int MAX_NNEI); -} // namespace deepmd diff --git a/source/op/tf/CMakeLists.txt b/source/op/tf/CMakeLists.txt index 6fc6422c6e..621037df85 100644 --- a/source/op/tf/CMakeLists.txt +++ b/source/op/tf/CMakeLists.txt @@ -34,7 +34,6 @@ file( ewald_recp.cc gelu_multi_device.cc map_aparam.cc - neighbor_stat.cc unaggregated_grad.cc tabulate_multi_device.cc prod_env_mat_multi_device.cc diff --git a/source/op/tf/neighbor_stat.cc b/source/op/tf/neighbor_stat.cc deleted file mode 100644 index 831c457a1c..0000000000 --- a/source/op/tf/neighbor_stat.cc +++ /dev/null @@ -1,329 +0,0 @@ -// SPDX-License-Identifier: LGPL-3.0-or-later -#include "neighbor_stat.h" - -#include "custom_op.h" -#include "errors.h" -#include "neighbor_list.h" - -typedef double boxtensor_t; -typedef double compute_t; - -REGISTER_OP("NeighborStat") - .Attr("T: {float, double} = DT_DOUBLE") - .Input("coord: T") - .Input("type: int32") - .Input("natoms: int32") - .Input("box : T") - .Input("mesh : int32") - .Attr("rcut: float") - .Output("max_nbor_size: int32") - .Output("min_nbor_dist: T"); - -template -class NeighborStatOp : public OpKernel { - public: - explicit NeighborStatOp(OpKernelConstruction* context) : OpKernel(context) { - OP_REQUIRES_OK(context, context->GetAttr("rcut", &rcut)); - max_nbor_size_nlist = 1024; - max_cpy_trial = 100; - mem_cpy = 256; - max_nnei_trial = 100; - mem_nnei = 256; - } - - void Compute(OpKernelContext* context) override { - deepmd::safe_compute( - context, [this](OpKernelContext* context) { this->_Compute(context); }); - } - - void _Compute(OpKernelContext* context) { - // Grab the input tensor - int context_input_index = 0; - const Tensor& coord_tensor = context->input(context_input_index++); - const Tensor& type_tensor = context->input(context_input_index++); - const Tensor& natoms_tensor = context->input(context_input_index++); - const Tensor& box_tensor = context->input(context_input_index++); - const Tensor& mesh_tensor = context->input(context_input_index++); - - OP_REQUIRES(context, (coord_tensor.shape().dims() == 2), - deepmd::tf_compat::InvalidArgument("Dim of coord should be 2")); - OP_REQUIRES(context, (type_tensor.shape().dims() == 2), - deepmd::tf_compat::InvalidArgument("Dim of type should be 2")); - OP_REQUIRES( - context, (natoms_tensor.shape().dims() == 1), - deepmd::tf_compat::InvalidArgument("Dim of natoms should be 1")); - OP_REQUIRES(context, (box_tensor.shape().dims() == 2), - deepmd::tf_compat::InvalidArgument("Dim of box should be 2")); - OP_REQUIRES(context, (mesh_tensor.shape().dims() == 1), - deepmd::tf_compat::InvalidArgument("Dim of mesh should be 1")); - OP_REQUIRES(context, (natoms_tensor.shape().dim_size(0) >= 3), - deepmd::tf_compat::InvalidArgument( - "number of atoms should be larger than (or equal to) 3")); - int nloc = natoms_tensor.flat().data()[0]; - int nall = natoms_tensor.flat().data()[1]; - int nsamples = coord_tensor.shape().dim_size(0); - int ntypes = natoms_tensor.shape().dim_size(0) - 2; - // check the sizes - OP_REQUIRES( - context, (nsamples == type_tensor.shape().dim_size(0)), - deepmd::tf_compat::InvalidArgument("number of samples should match")); - OP_REQUIRES( - context, (nsamples == box_tensor.shape().dim_size(0)), - deepmd::tf_compat::InvalidArgument("number of samples should match")); - OP_REQUIRES( - context, (nall * 3 == coord_tensor.shape().dim_size(1)), - deepmd::tf_compat::InvalidArgument("number of atoms should match")); - OP_REQUIRES( - context, (nall == type_tensor.shape().dim_size(1)), - deepmd::tf_compat::InvalidArgument("number of atoms should match")); - OP_REQUIRES( - context, (9 == box_tensor.shape().dim_size(1)), - deepmd::tf_compat::InvalidArgument("number of box should be 9")); - DeviceFunctor()(device, context->eigen_device()); - int nei_mode = 0; - if (mesh_tensor.shape().dim_size(0) == 6 || - mesh_tensor.shape().dim_size(0) == 7) { - // manual copied pbc - assert(nloc == nall); - nei_mode = 1; - } else if (mesh_tensor.shape().dim_size(0) == 0 || - mesh_tensor.shape().dim_size(0) == 1) { - // no pbc - nei_mode = -1; - } else { - throw deepmd::deepmd_exception("invalid mesh tensor"); - } - // if region is given extended, do not use pbc - bool b_pbc = (nei_mode >= 1 || nei_mode == -1) ? false : true; - bool b_norm_atom = (nei_mode == 1) ? true : false; - - TensorShape max_nbor_size_shape; - max_nbor_size_shape.AddDim(nloc); - max_nbor_size_shape.AddDim(ntypes); - - int context_output_index = 0; - Tensor* max_nbor_size_tensor = NULL; - OP_REQUIRES_OK(context, context->allocate_output(context_output_index++, - max_nbor_size_shape, - &max_nbor_size_tensor)); - - const FPTYPE* coord = coord_tensor.flat().data(); - const int* type = type_tensor.flat().data(); - const FPTYPE* box = box_tensor.flat().data(); - const int* mesh = mesh_tensor.flat().data(); - int* max_nbor_size = max_nbor_size_tensor->flat().data(); - if (device == "GPU") { -#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM - std::vector tensor_list(7); - if (nei_mode == 1) { - // Tensor FPTYPE_temp; - TensorShape FPTYPE_shape; - FPTYPE_shape.AddDim(static_cast(nall) * 3); - OP_REQUIRES_OK(context, - context->allocate_temp(DataTypeToEnum::value, - FPTYPE_shape, &tensor_list[0])); - - // Tensor double_temp; - TensorShape double_shape; - double_shape.AddDim(18); - OP_REQUIRES_OK(context, - context->allocate_temp(DataTypeToEnum::value, - double_shape, &tensor_list[1])); - // Tensor cpy_temp; - TensorShape cpy_shape; - cpy_shape.AddDim(static_cast(mem_cpy) * 3); - OP_REQUIRES_OK(context, - context->allocate_temp(DataTypeToEnum::value, - cpy_shape, &tensor_list[3])); - // Tensor t_temp; - TensorShape t_shape; - t_shape.AddDim(static_cast(mem_cpy) * 2); - OP_REQUIRES_OK(context, context->allocate_temp(DT_INT32, t_shape, - &tensor_list[4])); - } - - // Tensor nlist_temp; - TensorShape nlist_shape; - nlist_shape.AddDim(static_cast(nloc) * 2); - OP_REQUIRES_OK(context, context->allocate_temp(DT_INT32, nlist_shape, - &tensor_list[5])); - - TensorShape jlist_shape; - jlist_shape.AddDim(3 * int_64(nloc) * mem_nnei); - OP_REQUIRES_OK(context, context->allocate_temp(DT_INT32, jlist_shape, - &tensor_list[6])); - - int* idx_mapping = NULL; - int *ilist = NULL, *numneigh = NULL; - struct FirstneighGuard { - int** ptr = NULL; - ~FirstneighGuard() { - if (ptr != NULL) { - deepmd::delete_device_memory(ptr); - } - } - } firstneigh_guard; - deepmd::malloc_device_memory(firstneigh_guard.ptr, nloc); - int** firstneigh = firstneigh_guard.ptr; - int* jlist = NULL; - FPTYPE* coord_cpy; - int* type_cpy; - int frame_nall = nall; - int mesh_tensor_size = static_cast(mesh_tensor.NumElements()); - deepmd::InputNlist gpu_inlist; - int* nbor_list_dev = NULL; - // prepare coord and nlist - OP_REQUIRES_OK( - context, - _prepare_coord_nlist_gpu( - context, &tensor_list[0], &coord, coord_cpy, &type, type_cpy, - idx_mapping, gpu_inlist, ilist, numneigh, firstneigh, jlist, - nbor_list_dev, frame_nall, mem_cpy, mem_nnei, max_nbor_size_nlist, - box, mesh_tensor.flat().data(), mesh_tensor_size, nloc, - nei_mode, rcut, max_cpy_trial, max_nnei_trial)); - - TensorShape min_nbor_dist_shape; - min_nbor_dist_shape.AddDim(static_cast(nloc) * mem_nnei); - Tensor* min_nbor_dist_tensor = NULL; - OP_REQUIRES_OK(context, context->allocate_output(context_output_index++, - min_nbor_dist_shape, - &min_nbor_dist_tensor)); - FPTYPE* min_nbor_dist = min_nbor_dist_tensor->flat().data(); - - deepmd::neighbor_stat_gpu(coord, type, nloc, gpu_inlist, - max_nbor_size, min_nbor_dist, ntypes, - mem_nnei); -#endif - } else { - for (int ii = 0; - ii < static_cast(max_nbor_size_tensor->NumElements()); ii++) { - max_nbor_size[ii] = 0; - } - - // set region - boxtensor_t boxt[9] = {0}; - for (int dd = 0; dd < 9; ++dd) { - boxt[dd] = box[dd]; - } - SimulationRegion region; - region.reinitBox(boxt); - // set & normalize coord - std::vector d_coord3(nall * 3); - for (int ii = 0; ii < nall; ++ii) { - for (int dd = 0; dd < 3; ++dd) { - d_coord3[ii * 3 + dd] = coord[ii * 3 + dd]; - } - if (b_norm_atom) { - compute_t inter[3]; - region.phys2Inter(inter, &d_coord3[3 * ii]); - for (int dd = 0; dd < 3; ++dd) { - if (inter[dd] < 0) { - inter[dd] += 1.; - } else if (inter[dd] >= 1) { - inter[dd] -= 1.; - } - } - region.inter2Phys(&d_coord3[3 * ii], inter); - } - } - - // set type - std::vector d_type(nall); - for (int ii = 0; ii < nall; ++ii) { - d_type[ii] = type[ii]; - } - - // build nlist - std::vector > d_nlist_a; - std::vector > d_nlist_r; - std::vector nlist_map; - bool b_nlist_map = false; - - if (nei_mode == 1) { - // std::cout << "I'm in nei_mode 1" << std::endl; - std::vector bk_d_coord3 = d_coord3; - std::vector bk_d_type = d_type; - std::vector ncell, ngcell; - copy_coord(d_coord3, d_type, nlist_map, ncell, ngcell, bk_d_coord3, - bk_d_type, rcut, region); - b_nlist_map = true; - std::vector nat_stt(3, 0); - std::vector ext_stt(3), ext_end(3); - for (int dd = 0; dd < 3; ++dd) { - ext_stt[dd] = -ngcell[dd]; - ext_end[dd] = ncell[dd] + ngcell[dd]; - } - ::build_nlist(d_nlist_a, d_nlist_r, d_coord3, nloc, -1, rcut, nat_stt, - ncell, ext_stt, ext_end, region, ncell); - } else if (nei_mode == -1) { - ::build_nlist(d_nlist_a, d_nlist_r, d_coord3, -1, rcut, NULL); - } else { - throw deepmd::deepmd_exception("unknown neighbor mode"); - } - - int MAX_NNEI = 0; - for (int ii = 0; ii < nloc; ii++) { - MAX_NNEI = - MAX_NNEI < d_nlist_r[ii].size() ? d_nlist_r[ii].size() : MAX_NNEI; - } - // allocate output tensor for deepmd-kit - TensorShape min_nbor_dist_shape; - min_nbor_dist_shape.AddDim(static_cast(nloc) * MAX_NNEI); - Tensor* min_nbor_dist_tensor = NULL; - OP_REQUIRES_OK(context, context->allocate_output(context_output_index++, - min_nbor_dist_shape, - &min_nbor_dist_tensor)); - FPTYPE* min_nbor_dist = min_nbor_dist_tensor->flat().data(); - for (int ii = 0; - ii < static_cast(min_nbor_dist_tensor->NumElements()); ii++) { - min_nbor_dist[ii] = 10000.0; - } - -#pragma omp parallel for - for (int ii = 0; ii < nloc; ii++) { - if (d_type[ii] < 0) { - continue; // virtual atom - } - for (int jj = 0; jj < d_nlist_r[ii].size(); jj++) { - int type = d_type[d_nlist_r[ii][jj]]; - if (type < 0) { - continue; // virtual atom - } - max_nbor_size[ii * ntypes + type] += 1; - compute_t rij[3] = { - d_coord3[d_nlist_r[ii][jj] * 3 + 0] - d_coord3[ii * 3 + 0], - d_coord3[d_nlist_r[ii][jj] * 3 + 1] - d_coord3[ii * 3 + 1], - d_coord3[d_nlist_r[ii][jj] * 3 + 2] - d_coord3[ii * 3 + 2]}; - // we do not need to do slow sqrt for every dist; instead do sqrt in - // the final step - min_nbor_dist[ii * MAX_NNEI + jj] = - rij[0] * rij[0] + rij[1] * rij[1] + rij[2] * rij[2]; - } - } - } - } - - private: - int nnei; - float rcut; - std::string device; - int max_nbor_size_nlist, max_cpy_trial, mem_cpy, max_nnei_trial, mem_nnei; -}; - -#define REGISTER_CPU(T) \ - REGISTER_KERNEL_BUILDER( \ - Name("NeighborStat").Device(DEVICE_CPU).TypeConstraint("T"), \ - NeighborStatOp); -REGISTER_CPU(float); -REGISTER_CPU(double); -#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM -#define REGISTER_GPU(T) \ - REGISTER_KERNEL_BUILDER(Name("NeighborStat") \ - .Device(DEVICE_GPU) \ - .TypeConstraint("T") \ - .HostMemory("natoms") \ - .HostMemory("box"), \ - NeighborStatOp); -REGISTER_GPU(float); -REGISTER_GPU(double); -#endif // GOOGLE_CUDA || TENSORFLOW_USE_ROCM