From 58ae124e52f30ae69d2b47d2f08f3a3aab574502 Mon Sep 17 00:00:00 2001 From: dzzz2001 Date: Tue, 12 May 2026 16:48:26 +0800 Subject: [PATCH 1/2] [Refactor] Shorten verbose function names in module_gint Drop redundant verb prefixes and noise tokens: transfer_hr_gint_to_hR -> hr_gint_to_hR, transfer_dm_2d_to_gint -> dm_2d_to_gint, *_sparseMatrix -> *_sparse, init_atom_pair_start_end_idx_ -> init_atom_pair_idx_, get_max_atoms_num_per_bgrid -> get_max_atoms_per_bgrid. Pure rename, no behavior change. Co-Authored-By: Claude Opus 4.7 (1M context) --- .../source_lcao/module_gint/batch_biggrid.h | 2 +- .../source_lcao/module_gint/gint_common.cpp | 28 +++++++++---------- source/source_lcao/module_gint/gint_common.h | 4 +-- .../source_lcao/module_gint/gint_dvlocal.cpp | 16 +++++------ source/source_lcao/module_gint/gint_dvlocal.h | 4 +-- source/source_lcao/module_gint/gint_fvl.cpp | 2 +- .../source_lcao/module_gint/gint_fvl_gpu.cpp | 2 +- .../source_lcao/module_gint/gint_fvl_meta.cpp | 2 +- .../module_gint/gint_fvl_meta_gpu.cpp | 2 +- .../module_gint/gint_interface.cpp | 4 +-- .../source_lcao/module_gint/gint_interface.h | 2 +- source/source_lcao/module_gint/gint_rho.cpp | 2 +- .../source_lcao/module_gint/gint_rho_gpu.cpp | 2 +- source/source_lcao/module_gint/gint_tau.cpp | 2 +- .../source_lcao/module_gint/gint_tau_gpu.cpp | 2 +- source/source_lcao/module_gint/gint_vl.cpp | 4 +-- .../source_lcao/module_gint/gint_vl_gpu.cpp | 4 +-- .../module_gint/gint_vl_metagga.cpp | 2 +- .../module_gint/gint_vl_metagga_gpu.cpp | 2 +- .../module_gint/kernel/phi_operator_gpu.cu | 2 +- .../source_lcao/module_gint/phi_operator.cpp | 4 +-- source/source_lcao/module_gint/phi_operator.h | 2 +- .../module_gint/test/test_gint_common.cpp | 2 +- source/source_lcao/spar_dh.cpp | 2 +- 24 files changed, 50 insertions(+), 50 deletions(-) diff --git a/source/source_lcao/module_gint/batch_biggrid.h b/source/source_lcao/module_gint/batch_biggrid.h index d4de77d1dbe..87b5af7cc7c 100644 --- a/source/source_lcao/module_gint/batch_biggrid.h +++ b/source/source_lcao/module_gint/batch_biggrid.h @@ -16,7 +16,7 @@ class BatchBigGrid int get_batch_size() const { return biggrids_.size(); } int get_atoms_num() const { return atoms_num_; } int get_phi_len() const { return phi_len_;} - int get_max_atoms_num_per_bgrid() const { return max_atoms_num_per_bgrid_; } + int get_max_atoms_per_bgrid() const { return max_atoms_num_per_bgrid_; } bool empty() {return atoms_num_ == 0; } static int get_max_batch_size() { return max_batch_size_; } static int get_max_atoms_num() { return max_atoms_num_; } diff --git a/source/source_lcao/module_gint/gint_common.cpp b/source/source_lcao/module_gint/gint_common.cpp index 7d3e287bb84..987fcbe5a0e 100644 --- a/source/source_lcao/module_gint/gint_common.cpp +++ b/source/source_lcao/module_gint/gint_common.cpp @@ -113,10 +113,10 @@ void compose_hr_gint(HContainer& hr_gint) } template -void transfer_hr_gint_to_hR(const HContainer& hr_gint, HContainer& hR) +void hr_gint_to_hR(const HContainer& hr_gint, HContainer& hR) { - ModuleBase::TITLE("Gint", "transfer_hr_gint_to_hR"); - ModuleBase::timer::start("Gint", "transfer_hr_gint_to_hR"); + ModuleBase::TITLE("Gint", "hr_gint_to_hR"); + ModuleBase::timer::start("Gint", "hr_gint_to_hR"); #ifdef __MPI int size = 0; MPI_Comm_size(MPI_COMM_WORLD, &size); @@ -131,7 +131,7 @@ void transfer_hr_gint_to_hR(const HContainer& hr_gint, HContainer& hR) #else hR.add(hr_gint); #endif - ModuleBase::timer::end("Gint", "transfer_hr_gint_to_hR"); + ModuleBase::timer::end("Gint", "hr_gint_to_hR"); } @@ -258,7 +258,7 @@ void merge_hr_part_to_hR(const std::vector>& hr_gint_ -// C++11-compatible helpers for transfer_dm_2d_to_gint: +// C++11-compatible helpers for dm_2d_to_gint: // SFINAE (enable_if) to select same-type vs cross-type code paths at compile time. // Same-type path (TGint == TDM): transfer directly @@ -294,13 +294,13 @@ gather_dm(const HContainer& dm_src, HContainer& dm_dst, // gint_info should not have been a parameter, but it was added to initialize dm_gint_full // In the future, we might try to remove the gint_info parameter template -void transfer_dm_2d_to_gint( +void dm_2d_to_gint( const GintInfo& gint_info, const std::vector*>& dm, std::vector>& dm_gint) { - ModuleBase::TITLE("Gint", "transfer_dm_2d_to_gint"); - ModuleBase::timer::start("Gint", "transfer_dm_2d_to_gint"); + ModuleBase::TITLE("Gint", "dm_2d_to_gint"); + ModuleBase::timer::start("Gint", "dm_2d_to_gint"); if (PARAM.inp.nspin != 4) { @@ -368,7 +368,7 @@ void transfer_dm_2d_to_gint( delete dm2d_tmp; #endif } - ModuleBase::timer::end("Gint", "transfer_dm_2d_to_gint"); + ModuleBase::timer::end("Gint", "dm_2d_to_gint"); } int globalIndex(int localindex, int nblk, int nprocs, int myproc) @@ -480,10 +480,10 @@ void wfc_2d_to_gint(const T* wfc_2d, template void compose_hr_gint(HContainer& hr_gint); template void compose_hr_gint(HContainer& hr_gint); -template void transfer_hr_gint_to_hR( +template void hr_gint_to_hR( const HContainer& hr_gint, HContainer& hR); -template void transfer_hr_gint_to_hR( +template void hr_gint_to_hR( const HContainer>& hr_gint, HContainer>& hR); template void cast_hcontainer_values( @@ -494,15 +494,15 @@ template void cast_hcontainer_values( HContainer& dst); template HContainer make_cast_hcontainer(const HContainer& src); template HContainer make_cast_hcontainer(const HContainer& src); -template void transfer_dm_2d_to_gint( +template void dm_2d_to_gint( const GintInfo& gint_info, const std::vector*>& dm, std::vector>& dm_gint); -template void transfer_dm_2d_to_gint( +template void dm_2d_to_gint( const GintInfo& gint_info, const std::vector*>& dm, std::vector>& dm_gint); -template void transfer_dm_2d_to_gint( +template void dm_2d_to_gint( const GintInfo& gint_info, const std::vector>*>& dm, std::vector>>& dm_gint); diff --git a/source/source_lcao/module_gint/gint_common.h b/source/source_lcao/module_gint/gint_common.h index e0348e0f7af..b859022e37f 100644 --- a/source/source_lcao/module_gint/gint_common.h +++ b/source/source_lcao/module_gint/gint_common.h @@ -16,14 +16,14 @@ namespace ModuleGint template - void transfer_hr_gint_to_hR(const HContainer& hr_gint, HContainer& hR); + void hr_gint_to_hR(const HContainer& hr_gint, HContainer& hR); // for nspin=4 case void merge_hr_part_to_hR(const std::vector>& hr_gint_tmp , hamilt::HContainer>* hR, const GintInfo& gint_info); template - void transfer_dm_2d_to_gint( + void dm_2d_to_gint( const GintInfo& gint_info, const std::vector*>& dm, std::vector>& dm_gint); diff --git a/source/source_lcao/module_gint/gint_dvlocal.cpp b/source/source_lcao/module_gint/gint_dvlocal.cpp index e47e8e0228f..401a8cbade3 100644 --- a/source/source_lcao/module_gint/gint_dvlocal.cpp +++ b/source/source_lcao/module_gint/gint_dvlocal.cpp @@ -56,7 +56,7 @@ void Gint_dvlocal::cal_hr_gint_() } } -void Gint_dvlocal::cal_dvlocal_R_sparseMatrix( +void Gint_dvlocal::cal_dvlocal_R_sparse( const int nspin, const int cspin, const int nlocal, @@ -66,8 +66,8 @@ void Gint_dvlocal::cal_dvlocal_R_sparseMatrix( const Grid_Driver& gdriver, LCAO_HS_Arrays& hs_arrays) { - ModuleBase::TITLE("Gint", "cal_dvlocal_R_sparseMatrix"); - ModuleBase::timer::start("Gint", "cal_dvlocal_R_sparseMatrix"); + ModuleBase::TITLE("Gint", "cal_dvlocal_R_sparse"); + ModuleBase::timer::start("Gint", "cal_dvlocal_R_sparse"); std::map, std::map>> pvdpRx_sparseMatrix; std::map, std::map>> pvdpRy_sparseMatrix; std::map, std::map>> pvdpRz_sparseMatrix; @@ -123,14 +123,14 @@ void Gint_dvlocal::cal_dvlocal_R_sparseMatrix( } } } - distribute_pvdpR_sparseMatrix(cspin, 0, nlocal, sparse_thr, pvdpRx_sparseMatrix, pv, hs_arrays); - distribute_pvdpR_sparseMatrix(cspin, 1, nlocal, sparse_thr, pvdpRy_sparseMatrix, pv, hs_arrays); - distribute_pvdpR_sparseMatrix(cspin, 2, nlocal, sparse_thr, pvdpRz_sparseMatrix, pv, hs_arrays); - ModuleBase::timer::end("Gint", "cal_dvlocal_R_sparseMatrix"); + distribute_pvdpR_sparse(cspin, 0, nlocal, sparse_thr, pvdpRx_sparseMatrix, pv, hs_arrays); + distribute_pvdpR_sparse(cspin, 1, nlocal, sparse_thr, pvdpRy_sparseMatrix, pv, hs_arrays); + distribute_pvdpR_sparse(cspin, 2, nlocal, sparse_thr, pvdpRz_sparseMatrix, pv, hs_arrays); + ModuleBase::timer::end("Gint", "cal_dvlocal_R_sparse"); } -void Gint_dvlocal::distribute_pvdpR_sparseMatrix( +void Gint_dvlocal::distribute_pvdpR_sparse( const int cspin, const int dim, const int nlocal, diff --git a/source/source_lcao/module_gint/gint_dvlocal.h b/source/source_lcao/module_gint/gint_dvlocal.h index 606efbac36e..e4a9abf9609 100644 --- a/source/source_lcao/module_gint/gint_dvlocal.h +++ b/source/source_lcao/module_gint/gint_dvlocal.h @@ -24,7 +24,7 @@ class Gint_dvlocal : public Gint void cal_dvlocal(); - void cal_dvlocal_R_sparseMatrix( + void cal_dvlocal_R_sparse( const int nspin, const int cspin, const int nlocal, @@ -39,7 +39,7 @@ class Gint_dvlocal : public Gint void cal_hr_gint_(); - void distribute_pvdpR_sparseMatrix( + void distribute_pvdpR_sparse( const int cspin, const int dim, const int nlocal, diff --git a/source/source_lcao/module_gint/gint_fvl.cpp b/source/source_lcao/module_gint/gint_fvl.cpp index 752a11c4a11..fbf0b5d1f49 100644 --- a/source/source_lcao/module_gint/gint_fvl.cpp +++ b/source/source_lcao/module_gint/gint_fvl.cpp @@ -11,7 +11,7 @@ void Gint_fvl::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_fvl"); ModuleBase::timer::start("Gint", "cal_gint_fvl"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_fvl_svl_(); ModuleBase::timer::end("Gint", "cal_gint_fvl"); } diff --git a/source/source_lcao/module_gint/gint_fvl_gpu.cpp b/source/source_lcao/module_gint/gint_fvl_gpu.cpp index f1a721a1b8d..a40ccb9e9b7 100644 --- a/source/source_lcao/module_gint/gint_fvl_gpu.cpp +++ b/source/source_lcao/module_gint/gint_fvl_gpu.cpp @@ -13,7 +13,7 @@ void Gint_fvl_gpu::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_fvl"); ModuleBase::timer::start("Gint", "cal_gint_fvl"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_fvl_svl_(); ModuleBase::timer::end("Gint", "cal_gint_fvl"); } diff --git a/source/source_lcao/module_gint/gint_fvl_meta.cpp b/source/source_lcao/module_gint/gint_fvl_meta.cpp index b5d9a891bd9..f2cc3ce594f 100644 --- a/source/source_lcao/module_gint/gint_fvl_meta.cpp +++ b/source/source_lcao/module_gint/gint_fvl_meta.cpp @@ -11,7 +11,7 @@ void Gint_fvl_meta::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_fvl"); ModuleBase::timer::start("Gint", "cal_gint_fvl"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_fvl_svl_(); ModuleBase::timer::end("Gint", "cal_gint_fvl"); } diff --git a/source/source_lcao/module_gint/gint_fvl_meta_gpu.cpp b/source/source_lcao/module_gint/gint_fvl_meta_gpu.cpp index 6534e9169df..501f41eec8f 100644 --- a/source/source_lcao/module_gint/gint_fvl_meta_gpu.cpp +++ b/source/source_lcao/module_gint/gint_fvl_meta_gpu.cpp @@ -13,7 +13,7 @@ void Gint_fvl_meta_gpu::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_fvl"); ModuleBase::timer::start("Gint", "cal_gint_fvl"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_fvl_svl_(); ModuleBase::timer::end("Gint", "cal_gint_fvl"); } diff --git a/source/source_lcao/module_gint/gint_interface.cpp b/source/source_lcao/module_gint/gint_interface.cpp index 151f9ec475d..f8b896fe255 100644 --- a/source/source_lcao/module_gint/gint_interface.cpp +++ b/source/source_lcao/module_gint/gint_interface.cpp @@ -178,7 +178,7 @@ void cal_gint_fvl_meta( } } -void cal_dvlocal_R_sparseMatrix( +void cal_dvlocal_R_sparse( const int nspin, const int npol, const int current_spin, @@ -192,7 +192,7 @@ void cal_dvlocal_R_sparseMatrix( { Gint_dvlocal gint_dvlocal(vr_eff, nspin, npol); gint_dvlocal.cal_dvlocal(); - gint_dvlocal.cal_dvlocal_R_sparseMatrix( + gint_dvlocal.cal_dvlocal_R_sparse( nspin, current_spin, nlocal, sparse_thr, pv, ucell, gdriver, hs_arrays); } diff --git a/source/source_lcao/module_gint/gint_interface.h b/source/source_lcao/module_gint/gint_interface.h index 8fa40af4947..cb9112cf4f4 100644 --- a/source/source_lcao/module_gint/gint_interface.h +++ b/source/source_lcao/module_gint/gint_interface.h @@ -55,7 +55,7 @@ void cal_gint_fvl_meta( ModuleBase::matrix* fvl, ModuleBase::matrix* svl); -void cal_dvlocal_R_sparseMatrix( +void cal_dvlocal_R_sparse( const int nspin, const int npol, const int current_spin, diff --git a/source/source_lcao/module_gint/gint_rho.cpp b/source/source_lcao/module_gint/gint_rho.cpp index f9c99b2aff4..634659bc88b 100644 --- a/source/source_lcao/module_gint/gint_rho.cpp +++ b/source/source_lcao/module_gint/gint_rho.cpp @@ -36,7 +36,7 @@ void Gint_rho::cal_gint_impl_() { rho_data[is] = get_rho_data_(is, rho_cache); } - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec); cal_rho_(dm_gint_vec, rho_data); transfer_rho_cache_(rho_cache); } diff --git a/source/source_lcao/module_gint/gint_rho_gpu.cpp b/source/source_lcao/module_gint/gint_rho_gpu.cpp index 5dedac31621..dfa6fcc8606 100644 --- a/source/source_lcao/module_gint/gint_rho_gpu.cpp +++ b/source/source_lcao/module_gint/gint_rho_gpu.cpp @@ -38,7 +38,7 @@ void Gint_rho_gpu::cal_gint_impl_() } // 2. Transfer dm from 2D parallel distribution to gint serial distribution - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec); // 3. Transfer dm to GPU std::vector> dm_gint_d_vec(nspin_); diff --git a/source/source_lcao/module_gint/gint_tau.cpp b/source/source_lcao/module_gint/gint_tau.cpp index 139742a01c2..57534062836 100644 --- a/source/source_lcao/module_gint/gint_tau.cpp +++ b/source/source_lcao/module_gint/gint_tau.cpp @@ -11,7 +11,7 @@ void Gint_tau::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_tau"); ModuleBase::timer::start("Gint", "cal_gint_tau"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_tau_(); ModuleBase::timer::end("Gint", "cal_gint_tau"); } diff --git a/source/source_lcao/module_gint/gint_tau_gpu.cpp b/source/source_lcao/module_gint/gint_tau_gpu.cpp index fcc6017046f..dc70fdf4a97 100644 --- a/source/source_lcao/module_gint/gint_tau_gpu.cpp +++ b/source/source_lcao/module_gint/gint_tau_gpu.cpp @@ -13,7 +13,7 @@ void Gint_tau_gpu::cal_gint() ModuleBase::TITLE("Gint", "cal_gint_tau"); ModuleBase::timer::start("Gint", "cal_gint_tau"); init_dm_gint_(); - transfer_dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); + dm_2d_to_gint(*gint_info_, dm_vec_, dm_gint_vec_); cal_tau_(); ModuleBase::timer::end("Gint", "cal_gint_tau"); } diff --git a/source/source_lcao/module_gint/gint_vl.cpp b/source/source_lcao/module_gint/gint_vl.cpp index 83bdcfa56cc..2ff8647bafb 100644 --- a/source/source_lcao/module_gint/gint_vl.cpp +++ b/source/source_lcao/module_gint/gint_vl.cpp @@ -60,7 +60,7 @@ const Real* get_vr_eff_data_(const double* vr_eff, int local_mgrid_num, inline void finalize_hr_gint_(HContainer& hr_gint, HContainer* hR) { compose_hr_gint(hr_gint); - transfer_hr_gint_to_hR(hr_gint, *hR); + hr_gint_to_hR(hr_gint, *hR); } template @@ -68,7 +68,7 @@ void finalize_hr_gint_(HContainer& hr_gint, HContainer* hR) { HContainer hr_gint_dp = make_cast_hcontainer(hr_gint); compose_hr_gint(hr_gint_dp); - transfer_hr_gint_to_hR(hr_gint_dp, *hR); + hr_gint_to_hR(hr_gint_dp, *hR); } template diff --git a/source/source_lcao/module_gint/gint_vl_gpu.cpp b/source/source_lcao/module_gint/gint_vl_gpu.cpp index 0ae475582e0..ef2643f4a28 100644 --- a/source/source_lcao/module_gint/gint_vl_gpu.cpp +++ b/source/source_lcao/module_gint/gint_vl_gpu.cpp @@ -32,7 +32,7 @@ void Gint_vl_gpu::cal_gint() inline void finalize_hr_gint_gpu_(HContainer& hr_gint, HContainer* hR) { compose_hr_gint(hr_gint); - transfer_hr_gint_to_hR(hr_gint, *hR); + hr_gint_to_hR(hr_gint, *hR); } // Helper: finalize hr_gint (non-double path — cast to double first) @@ -41,7 +41,7 @@ void finalize_hr_gint_gpu_(HContainer& hr_gint, HContainer* hR) { HContainer hr_gint_dp = make_cast_hcontainer(hr_gint); compose_hr_gint(hr_gint_dp); - transfer_hr_gint_to_hR(hr_gint_dp, *hR); + hr_gint_to_hR(hr_gint_dp, *hR); } template diff --git a/source/source_lcao/module_gint/gint_vl_metagga.cpp b/source/source_lcao/module_gint/gint_vl_metagga.cpp index 9ff463b4dad..1128e33bf6f 100644 --- a/source/source_lcao/module_gint/gint_vl_metagga.cpp +++ b/source/source_lcao/module_gint/gint_vl_metagga.cpp @@ -13,7 +13,7 @@ void Gint_vl_metagga::cal_gint() init_hr_gint_(); cal_hr_gint_(); compose_hr_gint(hr_gint_); - transfer_hr_gint_to_hR(hr_gint_, *hR_); + hr_gint_to_hR(hr_gint_, *hR_); ModuleBase::timer::end("Gint", "cal_gint_vl"); } diff --git a/source/source_lcao/module_gint/gint_vl_metagga_gpu.cpp b/source/source_lcao/module_gint/gint_vl_metagga_gpu.cpp index 34a5387530c..5ddef7a7132 100644 --- a/source/source_lcao/module_gint/gint_vl_metagga_gpu.cpp +++ b/source/source_lcao/module_gint/gint_vl_metagga_gpu.cpp @@ -15,7 +15,7 @@ void Gint_vl_metagga_gpu::cal_gint() init_hr_gint_(); cal_hr_gint_(); compose_hr_gint(hr_gint_); - transfer_hr_gint_to_hR(hr_gint_, *hR_); + hr_gint_to_hR(hr_gint_, *hR_); ModuleBase::timer::end("Gint", "cal_gint_vl"); } diff --git a/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu b/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu index c338fcba671..47d5139db39 100644 --- a/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu +++ b/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu @@ -425,7 +425,7 @@ void PhiOperatorGpu::phi_dot_dphi( const double* dphi_z_d, double* fvl_d) const { - dim3 grid_dim(bgrid_batch_->get_max_atoms_num_per_bgrid(), + dim3 grid_dim(bgrid_batch_->get_max_atoms_per_bgrid(), bgrid_batch_->get_batch_size()); dim3 threads_per_block(32); phi_dot_dphi_kernel<<>>( diff --git a/source/source_lcao/module_gint/phi_operator.cpp b/source/source_lcao/module_gint/phi_operator.cpp index 5df52f94531..76dc5f8b4b2 100644 --- a/source/source_lcao/module_gint/phi_operator.cpp +++ b/source/source_lcao/module_gint/phi_operator.cpp @@ -29,7 +29,7 @@ void PhiOperator::set_bgrid(std::shared_ptr biggrid) } // init atom_pair_start_end_idx_ - init_atom_pair_start_end_idx_(); + init_atom_pair_idx_(); } void PhiOperator::set_phi_dphi(double* phi, double* dphi_x, double* dphi_y, double* dphi_z) const @@ -214,7 +214,7 @@ void PhiOperator::cal_env_k( //=============================== // private methods //=============================== -void PhiOperator::init_atom_pair_start_end_idx_() +void PhiOperator::init_atom_pair_idx_() { int atoms_num = biggrid_->get_atoms_num(); atom_pair_start_end_idx_.resize(atoms_num * (atoms_num + 1) / 2); diff --git a/source/source_lcao/module_gint/phi_operator.h b/source/source_lcao/module_gint/phi_operator.h index e5cdb168b50..323f99e8e93 100644 --- a/source/source_lcao/module_gint/phi_operator.h +++ b/source/source_lcao/module_gint/phi_operator.h @@ -112,7 +112,7 @@ class PhiOperator double* rho) const; private: - void init_atom_pair_start_end_idx_(); + void init_atom_pair_idx_(); // get the index of the first and the last meshgrid that both atom a and atom b affect // Note that atom_pair_start_end_idx_ only stores the cases where a <= b, so this function is needed to retrieve the value diff --git a/source/source_lcao/module_gint/test/test_gint_common.cpp b/source/source_lcao/module_gint/test/test_gint_common.cpp index b9d52f4c4ae..4d6a21607ca 100644 --- a/source/source_lcao/module_gint/test/test_gint_common.cpp +++ b/source/source_lcao/module_gint/test/test_gint_common.cpp @@ -105,7 +105,7 @@ TEST_F(GintCommonTest, TransferDm2dToGintSupportsDoubleToFloat) std::vector> dm_gint; dm_gint.push_back(gint_info_->get_hr()); - ModuleGint::transfer_dm_2d_to_gint(*gint_info_, dm_vec, dm_gint); + ModuleGint::dm_2d_to_gint(*gint_info_, dm_vec, dm_gint); ASSERT_EQ(dm_gint.size(), 1); EXPECT_EQ(dm.get_ijr_info(), dm_gint[0].get_ijr_info()); diff --git a/source/source_lcao/spar_dh.cpp b/source/source_lcao/spar_dh.cpp index 1848901eeaa..f1d8900c9b0 100644 --- a/source/source_lcao/spar_dh.cpp +++ b/source/source_lcao/spar_dh.cpp @@ -112,7 +112,7 @@ void sparse_format::cal_dH(const UnitCell& ucell, = v_eff.nc * v_eff.nr > 0 ? &(v_eff(current_spin, 0)) : nullptr; if (!PARAM.globalv.gamma_only_local) { - ModuleGint::cal_dvlocal_R_sparseMatrix( + ModuleGint::cal_dvlocal_R_sparse( PARAM.inp.nspin, PARAM.globalv.npol, current_spin, PARAM.globalv.nlocal, sparse_thr, vr_eff1, pv, ucell, grid, HS_Arrays); } From e3ef6f02bc60c1198c51f71bc1fdcba7f61d4446 Mon Sep 17 00:00:00 2001 From: dzzz2001 Date: Tue, 12 May 2026 23:32:33 +0800 Subject: [PATCH 2/2] [Refactor] Shorten verbose member names in module_gint Co-Authored-By: Claude Opus 4.7 (1M context) --- .../source_lcao/module_gint/batch_biggrid.cpp | 6 +- .../source_lcao/module_gint/batch_biggrid.h | 8 +- source/source_lcao/module_gint/big_grid.cpp | 8 +- .../source_lcao/module_gint/biggrid_info.cpp | 28 +++--- source/source_lcao/module_gint/biggrid_info.h | 18 ++-- .../source_lcao/module_gint/gint_dvlocal.cpp | 54 +++++----- source/source_lcao/module_gint/gint_dvlocal.h | 4 +- source/source_lcao/module_gint/gint_info.cpp | 18 ++-- source/source_lcao/module_gint/gint_vl.cpp | 2 +- .../module_gint/gint_vl_metagga.cpp | 8 +- .../module_gint/gint_vl_metagga_nspin4.cpp | 8 +- .../module_gint/gint_vl_nspin4.cpp | 2 +- .../module_gint/kernel/phi_operator_gpu.cu | 82 ++++++++-------- .../module_gint/kernel/phi_operator_gpu.h | 8 +- .../module_gint/kernel/phi_operator_kernel.cu | 98 +++++++++---------- .../kernel/phi_operator_kernel.cuh | 40 ++++---- .../source_lcao/module_gint/meshgrid_info.h | 38 +++---- .../source_lcao/module_gint/phi_operator.cpp | 28 +++--- source/source_lcao/module_gint/phi_operator.h | 16 +-- .../source_lcao/module_gint/phi_operator.hpp | 12 +-- 20 files changed, 243 insertions(+), 243 deletions(-) diff --git a/source/source_lcao/module_gint/batch_biggrid.cpp b/source/source_lcao/module_gint/batch_biggrid.cpp index 8372506e465..f7b9669f8fb 100644 --- a/source/source_lcao/module_gint/batch_biggrid.cpp +++ b/source/source_lcao/module_gint/batch_biggrid.cpp @@ -6,7 +6,7 @@ namespace ModuleGint int BatchBigGrid::max_batch_size_ = 0; int BatchBigGrid::max_atoms_num_ = 0; int BatchBigGrid::max_phi_len_ = 0; -int BatchBigGrid::max_atom_pairs_num_ = 0; +int BatchBigGrid::max_atom_pairs_ = 0; BatchBigGrid::BatchBigGrid(std::vector> biggrids) { @@ -19,14 +19,14 @@ BatchBigGrid::BatchBigGrid(std::vector> biggrids) { max_nw_ = std::max(max_nw_, atom->get_nw()); } - max_atoms_num_per_bgrid_ = std::max(max_atoms_num_per_bgrid_, biggrid->get_atoms_num()); + max_atoms_per_bgrid_ = std::max(max_atoms_per_bgrid_, biggrid->get_atoms_num()); atoms_num_ += biggrid->get_atoms_num(); atom_pairs_num += std::pow(biggrid->get_atoms_num(), 2); phi_len_ += biggrid->get_phi_len() * biggrid->get_mgrids_num(); } max_atoms_num_ = std::max(max_atoms_num_, atoms_num_); max_phi_len_ = std::max(max_phi_len_, phi_len_); - max_atom_pairs_num_ = std::max(max_atom_pairs_num_, atom_pairs_num); + max_atom_pairs_ = std::max(max_atom_pairs_, atom_pairs_num); } diff --git a/source/source_lcao/module_gint/batch_biggrid.h b/source/source_lcao/module_gint/batch_biggrid.h index 87b5af7cc7c..f5d9b8bd4d3 100644 --- a/source/source_lcao/module_gint/batch_biggrid.h +++ b/source/source_lcao/module_gint/batch_biggrid.h @@ -16,12 +16,12 @@ class BatchBigGrid int get_batch_size() const { return biggrids_.size(); } int get_atoms_num() const { return atoms_num_; } int get_phi_len() const { return phi_len_;} - int get_max_atoms_per_bgrid() const { return max_atoms_num_per_bgrid_; } + int get_max_atoms_per_bgrid() const { return max_atoms_per_bgrid_; } bool empty() {return atoms_num_ == 0; } static int get_max_batch_size() { return max_batch_size_; } static int get_max_atoms_num() { return max_atoms_num_; } static int get_max_phi_len() { return max_phi_len_; } - static int get_max_atom_pairs_num() { return max_atom_pairs_num_; } + static int get_max_atom_pairs_num() { return max_atom_pairs_; } static std::shared_ptr get_bgrid_info() { return BigGrid::get_bgrid_info(); } private: @@ -35,7 +35,7 @@ class BatchBigGrid int atoms_num_ = 0; // the max number of atoms of a single biggrid - int max_atoms_num_per_bgrid_ = 0; + int max_atoms_per_bgrid_ = 0; // the max number of biggrids of a biggrids batch static int max_batch_size_; @@ -44,7 +44,7 @@ class BatchBigGrid // the max number of total wavefunctions of a biggrids batch static int max_phi_len_; // the max number of atom pairs of a biggrids batch - static int max_atom_pairs_num_; + static int max_atom_pairs_; }; } \ No newline at end of file diff --git a/source/source_lcao/module_gint/big_grid.cpp b/source/source_lcao/module_gint/big_grid.cpp index e20a0fb50a6..66d96debcb7 100644 --- a/source/source_lcao/module_gint/big_grid.cpp +++ b/source/source_lcao/module_gint/big_grid.cpp @@ -45,10 +45,10 @@ void BigGrid::set_atoms_phi_len(std::vector& phi_len) const void BigGrid::set_mgrids_coord(std::vector& coord) const { coord.resize(biggrid_info_->get_mgrids_num()); - Vec3d this_bgrid_coord = localcell_info_->get_bgrid_global_coord_3D(idx_); + Vec3d bgrid_coord = localcell_info_->get_bgrid_global_coord_3D(idx_); for(int im = 0; im < biggrid_info_->get_mgrids_num(); ++im) { - coord[im] = biggrid_info_->get_mgrid_coord(im) + this_bgrid_coord; + coord[im] = biggrid_info_->get_mgrid_coord(im) + bgrid_coord; } } @@ -83,14 +83,14 @@ void BigGrid::set_atom_relative_coords(const Vec3i bgrid_idx, const Vec3d tau_in Vec3i this_bgrid_idx = localcell_info_->get_bgrid_global_idx_3D(idx_); // the relative coordinates of this big grid and the atom - Vec3d bgrid_relative_coord + Vec3d bgrid_rcoord = unitcell_info_->get_relative_coord(bgrid_idx, this_bgrid_idx) + tau_in_bgrid; atom_coord.resize(biggrid_info_->get_mgrids_num()); for(int im = 0; im < biggrid_info_->get_mgrids_num(); ++im) { const Vec3d& mgrid_coord = biggrid_info_->get_mgrid_coord(im); - atom_coord[im] = mgrid_coord - bgrid_relative_coord; + atom_coord[im] = mgrid_coord - bgrid_rcoord; } } diff --git a/source/source_lcao/module_gint/biggrid_info.cpp b/source/source_lcao/module_gint/biggrid_info.cpp index 52d1e9f3947..061563735e4 100644 --- a/source/source_lcao/module_gint/biggrid_info.cpp +++ b/source/source_lcao/module_gint/biggrid_info.cpp @@ -16,21 +16,21 @@ BigGridInfo::BigGridInfo( biggrid_vec3_(biggrid_vec3), nmx_(nmx), nmy_(nmy), nmz_(nmz), nmxyz_(nmx*nmy*nmz) { - // initialize the biggrid_latvec0_ - biggrid_latvec0_.e11 = biggrid_vec1_.x; - biggrid_latvec0_.e12 = biggrid_vec1_.y; - biggrid_latvec0_.e13 = biggrid_vec1_.z; + // initialize the bgrid_latvec0_ + bgrid_latvec0_.e11 = biggrid_vec1_.x; + bgrid_latvec0_.e12 = biggrid_vec1_.y; + bgrid_latvec0_.e13 = biggrid_vec1_.z; - biggrid_latvec0_.e21 = biggrid_vec2_.x; - biggrid_latvec0_.e22 = biggrid_vec2_.y; - biggrid_latvec0_.e23 = biggrid_vec2_.z; + bgrid_latvec0_.e21 = biggrid_vec2_.x; + bgrid_latvec0_.e22 = biggrid_vec2_.y; + bgrid_latvec0_.e23 = biggrid_vec2_.z; - biggrid_latvec0_.e31 = biggrid_vec3_.x; - biggrid_latvec0_.e32 = biggrid_vec3_.y; - biggrid_latvec0_.e33 = biggrid_vec3_.z; + bgrid_latvec0_.e31 = biggrid_vec3_.x; + bgrid_latvec0_.e32 = biggrid_vec3_.y; + bgrid_latvec0_.e33 = biggrid_vec3_.z; // initialize the GT matrix - biggrid_GT_ = biggrid_latvec0_.Inverse(); + biggrid_GT_ = bgrid_latvec0_.Inverse(); // initialize the meshgrid_info_ meshgrid_info_ = std::make_shared( @@ -38,11 +38,11 @@ BigGridInfo::BigGridInfo( biggrid_vec2_ / static_cast(nmy), biggrid_vec3_ / static_cast(nmz)); - // initialize the meshgrid_coords_ - meshgrid_coords_.resize(nmxyz_); + // initialize the mgrid_coords_ + mgrid_coords_.resize(nmxyz_); for(int index_1d = 0; index_1d < nmxyz_; index_1d++) { - meshgrid_coords_[index_1d] = + mgrid_coords_[index_1d] = meshgrid_info_->get_cartesian_coord(mgrid_idx_1Dto3D(index_1d)); } ModuleBase::Memory::record("BigGridInfo::meshgrid_coords", (long long)nmxyz_ * sizeof(Vec3d), true); diff --git a/source/source_lcao/module_gint/biggrid_info.h b/source/source_lcao/module_gint/biggrid_info.h index 65fd225240a..2f574a0188e 100644 --- a/source/source_lcao/module_gint/biggrid_info.h +++ b/source/source_lcao/module_gint/biggrid_info.h @@ -24,8 +24,8 @@ class BigGridInfo ~BigGridInfo(); - Vec3d get_cartesian_coord(const Vec3d& index_3d) const { return index_3d * biggrid_latvec0_; } - Vec3d get_cartesian_coord(const Vec3i& index_3d) const { return index_3d * biggrid_latvec0_; } + Vec3d get_cartesian_coord(const Vec3d& index_3d) const { return index_3d * bgrid_latvec0_; } + Vec3d get_cartesian_coord(const Vec3i& index_3d) const { return index_3d * bgrid_latvec0_; } Vec3d get_direct_coord(const Vec3d& cart_coord) const { return cart_coord * biggrid_GT_; } // Return the maximum number of big grids that can fit inside a sphere of radius r, @@ -38,8 +38,8 @@ class BigGridInfo int get_nmz() const { return nmz_; } int get_mgrids_num() const { return nmxyz_; } - const std::vector& get_mgrids_coord() const { return meshgrid_coords_; } - const Vec3d& get_mgrid_coord(int index_1d) const { return meshgrid_coords_[index_1d]; } + const std::vector& get_mgrids_coord() const { return mgrid_coords_; } + const Vec3d& get_mgrid_coord(int index_1d) const { return mgrid_coords_[index_1d]; } std::shared_ptr get_mgrid_info() const { return meshgrid_info_; } @@ -63,12 +63,12 @@ class BigGridInfo // used to convert the (i, j, k) index of the big grid to the Cartesian coordinate // if biggrid_vec1_ is row vector, - // then biggrid_latvec0_ = [biggrid_vec1_; biggrid_vec2_; biggrid_vec3_], - // (i, j, k) * biggrid_latvec0_ = (x, y, z) - Matrix3 biggrid_latvec0_; + // then bgrid_latvec0_ = [biggrid_vec1_; biggrid_vec2_; biggrid_vec3_], + // (i, j, k) * bgrid_latvec0_ = (x, y, z) + Matrix3 bgrid_latvec0_; // used to convert the Cartesian coordinate to the (i, j, k) index of the big grid - // biggrid_GT_ = biggrid_latvec0_.Inverse() + // biggrid_GT_ = bgrid_latvec0_.Inverse() // (x, y, z) * biggrid_GT_ = (i, j, k) Matrix3 biggrid_GT_; @@ -95,7 +95,7 @@ class BigGridInfo // store the relative Cartesian coordinates of all meshgrids in the biggrid // the size of vector is nbxyz_ - std::vector meshgrid_coords_; + std::vector mgrid_coords_; }; } // namespace ModuleGint \ No newline at end of file diff --git a/source/source_lcao/module_gint/gint_dvlocal.cpp b/source/source_lcao/module_gint/gint_dvlocal.cpp index 401a8cbade3..dccad124553 100644 --- a/source/source_lcao/module_gint/gint_dvlocal.cpp +++ b/source/source_lcao/module_gint/gint_dvlocal.cpp @@ -49,9 +49,9 @@ void Gint_dvlocal::cal_hr_gint_() dphi_z.resize(phi_len); phi_op.set_phi_dphi(phi.data(), dphi_x.data(), dphi_y.data(), dphi_z.data()); phi_op.phi_mul_vldr3(vr_eff_, dr3_, phi.data(), phi_vldr3.data()); - phi_op.phi_mul_phi(phi_vldr3.data(), dphi_x.data(), pvdpRx, PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(phi_vldr3.data(), dphi_y.data(), pvdpRy, PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(phi_vldr3.data(), dphi_z.data(), pvdpRz, PhiOperator::Triangular_Matrix::Upper); + phi_op.phi_mul_phi(phi_vldr3.data(), dphi_x.data(), pvdpRx, PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(phi_vldr3.data(), dphi_y.data(), pvdpRy, PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(phi_vldr3.data(), dphi_z.data(), pvdpRz, PhiOperator::TriPart::Upper); } } } @@ -68,11 +68,11 @@ void Gint_dvlocal::cal_dvlocal_R_sparse( { ModuleBase::TITLE("Gint", "cal_dvlocal_R_sparse"); ModuleBase::timer::start("Gint", "cal_dvlocal_R_sparse"); - std::map, std::map>> pvdpRx_sparseMatrix; - std::map, std::map>> pvdpRy_sparseMatrix; - std::map, std::map>> pvdpRz_sparseMatrix; + std::map, std::map>> pvdpRx_sparse; + std::map, std::map>> pvdpRy_sparse; + std::map, std::map>> pvdpRz_sparse; - double temp_value_double = 0.0; + double tmp_d = 0.0; Vec3d tau1, dtau; for (int iap = 0; iap < pvdpRx.size_atom_pairs(); iap++) @@ -107,25 +107,25 @@ void Gint_dvlocal::cal_dvlocal_R_sparse( double temp_value = p_pvdpRx[iw_nowg]; if (std::abs(temp_value) > sparse_thr) { - pvdpRx_sparseMatrix[dR][start1 + iw][start2 + iw2] = temp_value; + pvdpRx_sparse[dR][start1 + iw][start2 + iw2] = temp_value; } temp_value = p_pvdpRy[iw_nowg]; if (std::abs(temp_value) > sparse_thr) { - pvdpRy_sparseMatrix[dR][start1 + iw][start2 + iw2] = temp_value; + pvdpRy_sparse[dR][start1 + iw][start2 + iw2] = temp_value; } temp_value = p_pvdpRz[iw_nowg]; if (std::abs(temp_value) > sparse_thr) { - pvdpRz_sparseMatrix[dR][start1 + iw][start2 + iw2] = temp_value; + pvdpRz_sparse[dR][start1 + iw][start2 + iw2] = temp_value; } } } } } - distribute_pvdpR_sparse(cspin, 0, nlocal, sparse_thr, pvdpRx_sparseMatrix, pv, hs_arrays); - distribute_pvdpR_sparse(cspin, 1, nlocal, sparse_thr, pvdpRy_sparseMatrix, pv, hs_arrays); - distribute_pvdpR_sparse(cspin, 2, nlocal, sparse_thr, pvdpRz_sparseMatrix, pv, hs_arrays); + distribute_pvdpR_sparse(cspin, 0, nlocal, sparse_thr, pvdpRx_sparse, pv, hs_arrays); + distribute_pvdpR_sparse(cspin, 1, nlocal, sparse_thr, pvdpRy_sparse, pv, hs_arrays); + distribute_pvdpR_sparse(cspin, 2, nlocal, sparse_thr, pvdpRz_sparse, pv, hs_arrays); ModuleBase::timer::end("Gint", "cal_dvlocal_R_sparse"); } @@ -134,10 +134,10 @@ void Gint_dvlocal::distribute_pvdpR_sparse( const int cspin, const int dim, const int nlocal, - const double sparse_threshold, + const double sparse_thr, const std::map, std::map>>& - pvdpR_sparseMatrix, + pvdpR_sparse, const Parallel_Orbitals& pv, LCAO_HS_Arrays& hs_arrays) { @@ -147,8 +147,8 @@ void Gint_dvlocal::distribute_pvdpR_sparse( int count = 0; for (const auto& R_coor: hs_arrays.all_R_coor) { - auto iter = pvdpR_sparseMatrix.find(R_coor); - if (iter != pvdpR_sparseMatrix.end()) + auto iter = pvdpR_sparse.find(R_coor); + if (iter != pvdpR_sparse.end()) { for (auto& row_loop: iter->second) { @@ -158,8 +158,8 @@ void Gint_dvlocal::distribute_pvdpR_sparse( auto minus_R_coor = -1 * R_coor; - iter = pvdpR_sparseMatrix.find(minus_R_coor); - if (iter != pvdpR_sparseMatrix.end()) + iter = pvdpR_sparse.find(minus_R_coor); + if (iter != pvdpR_sparse.end()) { for (auto& row_loop: iter->second) { @@ -186,8 +186,8 @@ void Gint_dvlocal::distribute_pvdpR_sparse( { tmp.assign(tmp.size(), 0); - auto iter = pvdpR_sparseMatrix.find(R_coor); - if (iter != pvdpR_sparseMatrix.end()) + auto iter = pvdpR_sparse.find(R_coor); + if (iter != pvdpR_sparse.end()) { if (trace_lo[row] >= 0) @@ -203,8 +203,8 @@ void Gint_dvlocal::distribute_pvdpR_sparse( } } - auto minus_R_iter = pvdpR_sparseMatrix.find(minus_R_coor); - if (minus_R_iter != pvdpR_sparseMatrix.end()) + auto minus_R_iter = pvdpR_sparse.find(minus_R_coor); + if (minus_R_iter != pvdpR_sparse.end()) { for (int col = 0; col < row; ++col) { @@ -231,13 +231,13 @@ void Gint_dvlocal::distribute_pvdpR_sparse( { if (pv.global2local_col(col) >= 0) { - if (std::abs(tmp[col]) > sparse_threshold) + if (std::abs(tmp[col]) > sparse_thr) { if (dim == 0) { double& value = hs_arrays.dHRx_sparse[cspin][R_coor][row][col]; value += tmp[col]; - if (std::abs(value) <= sparse_threshold) + if (std::abs(value) <= sparse_thr) { hs_arrays.dHRx_sparse[cspin][R_coor][row].erase(col); } @@ -246,7 +246,7 @@ void Gint_dvlocal::distribute_pvdpR_sparse( { double& value = hs_arrays.dHRy_sparse[cspin][R_coor][row][col]; value += tmp[col]; - if (std::abs(value) <= sparse_threshold) + if (std::abs(value) <= sparse_thr) { hs_arrays.dHRy_sparse[cspin][R_coor][row].erase(col); } @@ -255,7 +255,7 @@ void Gint_dvlocal::distribute_pvdpR_sparse( { double& value = hs_arrays.dHRz_sparse[cspin][R_coor][row][col]; value += tmp[col]; - if (std::abs(value) <= sparse_threshold) + if (std::abs(value) <= sparse_thr) { hs_arrays.dHRz_sparse[cspin][R_coor][row].erase(col); } diff --git a/source/source_lcao/module_gint/gint_dvlocal.h b/source/source_lcao/module_gint/gint_dvlocal.h index e4a9abf9609..613976e678f 100644 --- a/source/source_lcao/module_gint/gint_dvlocal.h +++ b/source/source_lcao/module_gint/gint_dvlocal.h @@ -43,10 +43,10 @@ class Gint_dvlocal : public Gint const int cspin, const int dim, const int nlocal, - const double sparse_threshold, + const double sparse_thr, const std::map, std::map>>& - pvdpR_sparseMatrix, + pvdpR_sparse, const Parallel_Orbitals& pv, LCAO_HS_Arrays& HS_Arrays); diff --git a/source/source_lcao/module_gint/gint_info.cpp b/source/source_lcao/module_gint/gint_info.cpp index 4e11a62952d..c162befb149 100644 --- a/source/source_lcao/module_gint/gint_info.cpp +++ b/source/source_lcao/module_gint/gint_info.cpp @@ -112,18 +112,18 @@ void GintInfo::init_atoms_(int ntype, const Atom* atoms, const Numerical_Orbital { // get the extended biggrid idx of the affected biggrid const Vec3i ext_bgrid_idx(bgrid_x, bgrid_y, bgrid_z); - const Vec3i normal_bgrid_idx = unitcell_info_->map_ext_idx_to_ucell(ext_bgrid_idx); - if(localcell_info_->is_bgrid_in_lcell(normal_bgrid_idx) == false) + const Vec3i norm_bgrid_idx = unitcell_info_->map_ext_idx_to_ucell(ext_bgrid_idx); + if(localcell_info_->is_bgrid_in_lcell(norm_bgrid_idx) == false) { continue; } - const int bgrid_local_idx = localcell_info_->get_bgrid_local_idx_1D(normal_bgrid_idx); + const int bgrid_local_idx = localcell_info_->get_bgrid_local_idx_1D(norm_bgrid_idx); // get the unitcell idx of the big grid const Vec3i ucell_idx_bgrid = unitcell_info_->get_unitcell_idx(ext_bgrid_idx); // The index of the unitcell containing the biggrid relative to the unitcell containing the atom. - const Vec3i ucell_idx_relative = ucell_idx_bgrid - ucell_idx_atom; - auto it = r_to_atom.find(ucell_idx_relative); + const Vec3i ucell_idx_rel = ucell_idx_bgrid - ucell_idx_atom; + auto it = r_to_atom.find(ucell_idx_rel); // if the gint_atom is not in the map, // it means this is the first time we find this atom may affect some biggrids, // add it to the r_to_atom map @@ -132,12 +132,12 @@ void GintInfo::init_atoms_(int ntype, const Atom* atoms, const Numerical_Orbital Vec3i ext_atom_bgrid_idx(atom_bgrid_idx.x - ucell_idx_bgrid.x * unitcell_info_->get_nbx(), atom_bgrid_idx.y - ucell_idx_bgrid.y * unitcell_info_->get_nby(), atom_bgrid_idx.z - ucell_idx_bgrid.z * unitcell_info_->get_nbz()); - r_to_atom.insert(std::make_pair(ucell_idx_relative, - GintAtom(&atom, i, j, iat, ext_atom_bgrid_idx, ucell_idx_relative, tau_in_biggrid, orb, ucell_))); + r_to_atom.insert(std::make_pair(ucell_idx_rel, + GintAtom(&atom, i, j, iat, ext_atom_bgrid_idx, ucell_idx_rel, tau_in_biggrid, orb, ucell_))); } - if(biggrids_[bgrid_local_idx]->is_atom_on_bgrid(&r_to_atom.at(ucell_idx_relative))) + if(biggrids_[bgrid_local_idx]->is_atom_on_bgrid(&r_to_atom.at(ucell_idx_rel))) { - biggrids_[bgrid_local_idx]->add_atom(&r_to_atom.at(ucell_idx_relative)); + biggrids_[bgrid_local_idx]->add_atom(&r_to_atom.at(ucell_idx_rel)); is_atom_in_proc_[iat] = true; } } diff --git a/source/source_lcao/module_gint/gint_vl.cpp b/source/source_lcao/module_gint/gint_vl.cpp index 2ff8647bafb..410eeb65885 100644 --- a/source/source_lcao/module_gint/gint_vl.cpp +++ b/source/source_lcao/module_gint/gint_vl.cpp @@ -98,7 +98,7 @@ void Gint_vl::cal_gint_impl_() phi_vldr3.resize(phi_len); phi_op.set_phi(phi.data()); phi_op.phi_mul_vldr3(vr_eff, static_cast(dr3_), phi.data(), phi_vldr3.data()); - phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint, PhiOperator::Triangular_Matrix::Upper); + phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint, PhiOperator::TriPart::Upper); } } diff --git a/source/source_lcao/module_gint/gint_vl_metagga.cpp b/source/source_lcao/module_gint/gint_vl_metagga.cpp index 1128e33bf6f..48b85f00a2b 100644 --- a/source/source_lcao/module_gint/gint_vl_metagga.cpp +++ b/source/source_lcao/module_gint/gint_vl_metagga.cpp @@ -62,10 +62,10 @@ void Gint_vl_metagga::cal_hr_gint_() phi_op.phi_mul_vldr3(vofk_, dr3_, dphi_x.data(), dphi_x_vldr3.data()); phi_op.phi_mul_vldr3(vofk_, dr3_, dphi_y.data(), dphi_y_vldr3.data()); phi_op.phi_mul_vldr3(vofk_, dr3_, dphi_z.data(), dphi_z_vldr3.data()); - phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_, PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_x.data(), dphi_x_vldr3.data(), hr_gint_, PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_y.data(), dphi_y_vldr3.data(), hr_gint_, PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_z.data(), dphi_z_vldr3.data(), hr_gint_, PhiOperator::Triangular_Matrix::Upper); + phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_, PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_x.data(), dphi_x_vldr3.data(), hr_gint_, PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_y.data(), dphi_y_vldr3.data(), hr_gint_, PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_z.data(), dphi_z_vldr3.data(), hr_gint_, PhiOperator::TriPart::Upper); } } } diff --git a/source/source_lcao/module_gint/gint_vl_metagga_nspin4.cpp b/source/source_lcao/module_gint/gint_vl_metagga_nspin4.cpp index ce7e27b2f11..bd2631aabc4 100644 --- a/source/source_lcao/module_gint/gint_vl_metagga_nspin4.cpp +++ b/source/source_lcao/module_gint/gint_vl_metagga_nspin4.cpp @@ -65,10 +65,10 @@ void Gint_vl_metagga_nspin4::cal_hr_gint_() phi_op.phi_mul_vldr3(vofk_[is], dr3_, dphi_x.data(), dphi_x_vldr3.data()); phi_op.phi_mul_vldr3(vofk_[is], dr3_, dphi_y.data(), dphi_y_vldr3.data()); phi_op.phi_mul_vldr3(vofk_[is], dr3_, dphi_z.data(), dphi_z_vldr3.data()); - phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_part_[is], PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_x.data(), dphi_x_vldr3.data(), hr_gint_part_[is], PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_y.data(), dphi_y_vldr3.data(), hr_gint_part_[is], PhiOperator::Triangular_Matrix::Upper); - phi_op.phi_mul_phi(dphi_z.data(), dphi_z_vldr3.data(), hr_gint_part_[is], PhiOperator::Triangular_Matrix::Upper); + phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_part_[is], PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_x.data(), dphi_x_vldr3.data(), hr_gint_part_[is], PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_y.data(), dphi_y_vldr3.data(), hr_gint_part_[is], PhiOperator::TriPart::Upper); + phi_op.phi_mul_phi(dphi_z.data(), dphi_z_vldr3.data(), hr_gint_part_[is], PhiOperator::TriPart::Upper); } } } diff --git a/source/source_lcao/module_gint/gint_vl_nspin4.cpp b/source/source_lcao/module_gint/gint_vl_nspin4.cpp index 7da30838ce8..d2157b8a37f 100644 --- a/source/source_lcao/module_gint/gint_vl_nspin4.cpp +++ b/source/source_lcao/module_gint/gint_vl_nspin4.cpp @@ -49,7 +49,7 @@ void Gint_vl_nspin4::cal_hr_gint_() for(int is = 0; is < nspin_; is++) { phi_op.phi_mul_vldr3(vr_eff_[is], dr3_, phi.data(), phi_vldr3.data()); - phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_part_[is], PhiOperator::Triangular_Matrix::Upper); + phi_op.phi_mul_phi(phi.data(), phi_vldr3.data(), hr_gint_part_[is], PhiOperator::TriPart::Upper); } } } diff --git a/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu b/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu index 47d5139db39..29e27a852a7 100644 --- a/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu +++ b/source/source_lcao/module_gint/kernel/phi_operator_gpu.cu @@ -12,12 +12,12 @@ PhiOperatorGpu::PhiOperatorGpu(std::shared_ptr gint_gpu :gint_gpu_vars_(gint_gpu_vars), stream_(stream), mgrids_num_(BatchBigGrid::get_bgrid_info()->get_mgrids_num()), atoms_num_info_(BatchBigGrid::get_max_batch_size(), stream_, true), -bgrids_phi_len_(BatchBigGrid::get_max_batch_size(), stream_, true), -bgrids_phi_start_(BatchBigGrid::get_max_batch_size(), stream_, true), +bgrid_phi_len_(BatchBigGrid::get_max_batch_size(), stream_, true), +bgrid_phi_start_(BatchBigGrid::get_max_batch_size(), stream_, true), atoms_iat_(BatchBigGrid::get_max_atoms_num(), stream_, true), atoms_bgrids_rcoords_(BatchBigGrid::get_max_atoms_num(), stream_, true), -atoms_phi_start_(BatchBigGrid::get_max_atoms_num(), stream_, true), -mgrids_local_idx_batch_(BatchBigGrid::get_max_batch_size() +atom_phi_start_(BatchBigGrid::get_max_atoms_num(), stream_, true), +batch_mgrid_lidx_(BatchBigGrid::get_max_batch_size() * BatchBigGrid::get_bgrid_info()->get_mgrids_num(), stream_, true), gemm_m_(BatchBigGrid::get_max_atom_pairs_num(), stream_, true), gemm_n_(BatchBigGrid::get_max_atom_pairs_num(), stream_, true), @@ -44,50 +44,50 @@ void PhiOperatorGpu::set_bgrid_batch(std::shared_ptr bgrid_b { bgrid_batch_ = bgrid_batch; auto atoms_num_info_h = atoms_num_info_.get_host_ptr(); - auto bgrids_phi_len_h = bgrids_phi_len_.get_host_ptr(); - auto bgrids_phi_start_h = bgrids_phi_start_.get_host_ptr(); + auto bgrid_phi_len_h = bgrid_phi_len_.get_host_ptr(); + auto bgrid_phi_start_h = bgrid_phi_start_.get_host_ptr(); auto atoms_iat_h = atoms_iat_.get_host_ptr(); - auto atoms_bgrids_rcoords_h = atoms_bgrids_rcoords_.get_host_ptr(); - auto atoms_phi_start_h = atoms_phi_start_.get_host_ptr(); - auto mgrids_local_idx_batch_h = mgrids_local_idx_batch_.get_host_ptr(); + auto atom_rcoords_h = atoms_bgrids_rcoords_.get_host_ptr(); + auto atom_phi_start_h = atom_phi_start_.get_host_ptr(); + auto batch_mgrid_lidx_h = batch_mgrid_lidx_.get_host_ptr(); int i = 0; int j = 0; int atoms_accum = 0; phi_len_ = 0; int phi_start = 0; - std::vector mgrids_local_idx; + std::vector mgrid_lidx; CHECK_CUDA(cudaEventSynchronize(event_)); for (const auto& bgrid : bgrid_batch->get_bgrids()) { atoms_num_info_h[i] = make_int2(bgrid->get_atoms_num(), atoms_accum); atoms_accum += bgrid->get_atoms_num(); - bgrids_phi_start_h[i] = phi_start; - bgrid->set_mgrids_local_idx(mgrids_local_idx); - std::copy(mgrids_local_idx.begin(), mgrids_local_idx.end(), - mgrids_local_idx_batch_h + i * mgrids_num_); + bgrid_phi_start_h[i] = phi_start; + bgrid->set_mgrids_local_idx(mgrid_lidx); + std::copy(mgrid_lidx.begin(), mgrid_lidx.end(), + batch_mgrid_lidx_h + i * mgrids_num_); int phi_len_bgrid = 0; for (const auto& atom : bgrid->get_atoms()) { atoms_iat_h[j] = atom->get_iat(); Vec3d rcoord = bgrid->get_bgrid_atom_rcoord(atom); - atoms_bgrids_rcoords_h[j] = make_double3(rcoord.x, rcoord.y, rcoord.z); - atoms_phi_start_h[j] = phi_len_ + phi_len_bgrid; + atom_rcoords_h[j] = make_double3(rcoord.x, rcoord.y, rcoord.z); + atom_phi_start_h[j] = phi_len_ + phi_len_bgrid; phi_len_bgrid += atom->get_nw(); j++; } - bgrids_phi_len_h[i] = phi_len_bgrid; + bgrid_phi_len_h[i] = phi_len_bgrid; phi_len_ += phi_len_bgrid * bgrid->get_mgrids_num(); phi_start += phi_len_bgrid * bgrid->get_mgrids_num(); i++; } atoms_num_info_.copy_host_to_device_async(bgrid_batch->get_batch_size()); - bgrids_phi_len_.copy_host_to_device_async(bgrid_batch->get_batch_size()); - bgrids_phi_start_.copy_host_to_device_async(bgrid_batch->get_batch_size()); + bgrid_phi_len_.copy_host_to_device_async(bgrid_batch->get_batch_size()); + bgrid_phi_start_.copy_host_to_device_async(bgrid_batch->get_batch_size()); atoms_iat_.copy_host_to_device_async(bgrid_batch->get_atoms_num()); atoms_bgrids_rcoords_.copy_host_to_device_async(bgrid_batch->get_atoms_num()); - atoms_phi_start_.copy_host_to_device_async(bgrid_batch->get_atoms_num()); - mgrids_local_idx_batch_.copy_host_to_device_async(bgrid_batch->get_batch_size() * mgrids_num_); + atom_phi_start_.copy_host_to_device_async(bgrid_batch->get_atoms_num()); + batch_mgrid_lidx_.copy_host_to_device_async(bgrid_batch->get_batch_size() * mgrids_num_); CHECK_CUDA(cudaEventRecord(event_, stream_)); } @@ -113,8 +113,8 @@ void PhiOperatorGpu::set_phi(Real* phi_d) const atoms_iat_.get_device_ptr(), atoms_bgrids_rcoords_.get_device_ptr(), atoms_num_info_.get_device_ptr(), - atoms_phi_start_.get_device_ptr(), - bgrids_phi_len_.get_device_ptr(), + atom_phi_start_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), phi_d); CHECK_LAST_CUDA_ERROR("kernel launch"); } @@ -142,8 +142,8 @@ void PhiOperatorGpu::set_phi_dphi(double* phi_d, double* dphi_x_d, double* atoms_iat_.get_device_ptr(), atoms_bgrids_rcoords_.get_device_ptr(), atoms_num_info_.get_device_ptr(), - atoms_phi_start_.get_device_ptr(), - bgrids_phi_len_.get_device_ptr(), + atom_phi_start_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), phi_d, dphi_x_d, dphi_y_d, @@ -183,8 +183,8 @@ void PhiOperatorGpu::set_ddphi(double* ddphi_xx_d, double* ddphi_xy_d, dou atoms_iat_.get_device_ptr(), atoms_bgrids_rcoords_.get_device_ptr(), atoms_num_info_.get_device_ptr(), - atoms_phi_start_.get_device_ptr(), - bgrids_phi_len_.get_device_ptr(), + atom_phi_start_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), ddphi_xx_d, ddphi_xy_d, ddphi_xz_d, @@ -208,9 +208,9 @@ void PhiOperatorGpu::phi_mul_vldr3( dr3, phi_d, mgrids_num_, - mgrids_local_idx_batch_.get_device_ptr(), - bgrids_phi_len_.get_device_ptr(), - bgrids_phi_start_.get_device_ptr(), + batch_mgrid_lidx_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), + bgrid_phi_start_.get_device_ptr(), result_d); CHECK_LAST_CUDA_ERROR("kernel launch"); } @@ -240,7 +240,7 @@ void PhiOperatorGpu::phi_mul_phi( const int iat_1 = atom_1->get_iat(); const auto& r_1 = atom_1->get_R(); const int nw1 = atom_1->get_nw(); - const int phi_1_offset = atoms_phi_start_.get_host_ptr()[pre_atoms + ia_1]; + const int phi_1_offset = atom_phi_start_.get_host_ptr()[pre_atoms + ia_1]; for (int ia_2 = 0; ia_2 < bgrid->get_atoms_num(); ia_2++) { @@ -256,7 +256,7 @@ void PhiOperatorGpu::phi_mul_phi( if (hr_offset == -1) { continue; } - const int phi_2_offset = atoms_phi_start_.get_host_ptr()[pre_atoms + ia_2]; + const int phi_2_offset = atom_phi_start_.get_host_ptr()[pre_atoms + ia_2]; gemm_A_.get_host_ptr()[ap_num] = phi_d + phi_1_offset; gemm_B_.get_host_ptr()[ap_num] = phi_vldr3_d + phi_2_offset; @@ -330,7 +330,7 @@ void PhiOperatorGpu::phi_mul_dm( const int iat_1 = atom_1->get_iat(); const auto& r_1 = atom_1->get_R(); const int nw1 = atom_1->get_nw(); - const int phi_1_offset = atoms_phi_start_.get_host_ptr()[pre_atoms + ia_1]; + const int phi_1_offset = atom_phi_start_.get_host_ptr()[pre_atoms + ia_1]; int ia_2 = is_symm ? ia_1 : 0; for (; ia_2 < bgrid->get_atoms_num(); ia_2++) { @@ -343,7 +343,7 @@ void PhiOperatorGpu::phi_mul_dm( if (dm_offset == -1) { continue; } - const int phi_dm_offset = atoms_phi_start_.get_host_ptr()[pre_atoms + ia_2]; + const int phi_dm_offset = atom_phi_start_.get_host_ptr()[pre_atoms + ia_2]; gemm_A_.get_host_ptr()[ap_num] = phi_d + phi_1_offset; gemm_B_.get_host_ptr()[ap_num] = dm_d + dm_offset; @@ -410,9 +410,9 @@ void PhiOperatorGpu::phi_dot_phi( phi_i_d, phi_j_d, mgrids_num_, - mgrids_local_idx_batch_.get_device_ptr(), - bgrids_phi_len_.get_device_ptr(), - bgrids_phi_start_.get_device_ptr(), + batch_mgrid_lidx_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), + bgrid_phi_start_.get_device_ptr(), rho_d); CHECK_LAST_CUDA_ERROR("kernel launch"); } @@ -434,9 +434,9 @@ void PhiOperatorGpu::phi_dot_dphi( dphi_y_d, dphi_z_d, mgrids_num_, - bgrids_phi_len_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), atoms_num_info_.get_device_ptr(), - atoms_phi_start_.get_device_ptr(), + atom_phi_start_.get_device_ptr(), atoms_iat_.get_device_ptr(), gint_gpu_vars_->iat2it_d, gint_gpu_vars_->atom_nw_d, @@ -461,9 +461,9 @@ void PhiOperatorGpu::phi_dot_dphi_r( dphi_y_d, dphi_z_d, mgrids_num_, - bgrids_phi_len_.get_device_ptr(), + bgrid_phi_len_.get_device_ptr(), atoms_num_info_.get_device_ptr(), - atoms_phi_start_.get_device_ptr(), + atom_phi_start_.get_device_ptr(), atoms_iat_.get_device_ptr(), atoms_bgrids_rcoords_.get_device_ptr(), gint_gpu_vars_->mgrids_pos_d, diff --git a/source/source_lcao/module_gint/kernel/phi_operator_gpu.h b/source/source_lcao/module_gint/kernel/phi_operator_gpu.h index 188129db679..ed440f499d6 100644 --- a/source/source_lcao/module_gint/kernel/phi_operator_gpu.h +++ b/source/source_lcao/module_gint/kernel/phi_operator_gpu.h @@ -89,13 +89,13 @@ class PhiOperatorGpu CudaMemWrapper atoms_bgrids_rcoords_; // the start index of the phi array for each atom - CudaMemWrapper atoms_phi_start_; + CudaMemWrapper atom_phi_start_; // The length of phi for a single meshgrid on each big grid. - CudaMemWrapper bgrids_phi_len_; + CudaMemWrapper bgrid_phi_len_; // The start index of the phi array for each big grid. - CudaMemWrapper bgrids_phi_start_; + CudaMemWrapper bgrid_phi_start_; // Mapping of the index of meshgrid in the batch of biggrids to the index of meshgrid in the local cell - CudaMemWrapper mgrids_local_idx_batch_; + CudaMemWrapper batch_mgrid_lidx_; mutable CudaMemWrapper gemm_m_; mutable CudaMemWrapper gemm_n_; diff --git a/source/source_lcao/module_gint/kernel/phi_operator_kernel.cu b/source/source_lcao/module_gint/kernel/phi_operator_kernel.cu index 1bb686de960..aa2dcc9fb7d 100644 --- a/source/source_lcao/module_gint/kernel/phi_operator_kernel.cu +++ b/source/source_lcao/module_gint/kernel/phi_operator_kernel.cu @@ -23,10 +23,10 @@ __global__ void set_phi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, Real* __restrict__ phi) { const int bgrid_id = blockIdx.y; @@ -38,7 +38,7 @@ __global__ void set_phi_kernel( for (int atom_id = threadIdx.x; atom_id < atoms_num; atom_id += blockDim.x) { const int atom_type = iat2it[atoms_iat[atom_id + pre_atoms_num]]; - const double3 rcoord = atoms_bgrids_rcoords[atom_id + pre_atoms_num]; // rcoord is the ralative coordinate of an atom and a biggrid + const double3 rcoord = atom_rcoords[atom_id + pre_atoms_num]; // rcoord is the ralative coordinate of an atom and a biggrid const double3 coord = make_double3(mgrid_pos.x-rcoord.x, // coord is the relative coordinate of an atom and a meshgrid mgrid_pos.y-rcoord.y, mgrid_pos.z-rcoord.z); @@ -66,8 +66,8 @@ __global__ void set_phi_kernel( double psi = 0; const int it_nw = atom_type * nwmax; int iw_nr = it_nw * nrmax + ip; - int phi_idx = atoms_phi_start[atom_id + pre_atoms_num] + - bgrids_phi_len[bgrid_id] * mgrid_id; + int phi_idx = atom_phi_start[atom_id + pre_atoms_num] + + bgrid_phi_len[bgrid_id] * mgrid_id; for (int iw = 0; iw < atom_nw[atom_type]; iw++, iw_nr += nrmax) { @@ -81,8 +81,8 @@ __global__ void set_phi_kernel( } else { - int phi_idx = atoms_phi_start[atom_id + pre_atoms_num] + - bgrids_phi_len[bgrid_id] * mgrid_id; + int phi_idx = atom_phi_start[atom_id + pre_atoms_num] + + bgrid_phi_len[bgrid_id] * mgrid_id; for (int iw = 0; iw < atom_nw[atom_type]; iw++) { phi[phi_idx + iw] = Real(0.0); @@ -121,10 +121,10 @@ __global__ void set_phi_dphi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, double* __restrict__ phi, double* __restrict__ dphi_x, double* __restrict__ dphi_y, @@ -139,7 +139,7 @@ __global__ void set_phi_dphi_kernel( for (int atom_id = threadIdx.x; atom_id < atoms_num; atom_id += blockDim.x) { const int atom_type = iat2it[atoms_iat[atom_id + pre_atoms_num]]; - const double3 rcoord = atoms_bgrids_rcoords[atom_id + pre_atoms_num]; + const double3 rcoord = atom_rcoords[atom_id + pre_atoms_num]; const double3 coord = make_double3(mgrid_pos.x-rcoord.x, mgrid_pos.y-rcoord.y, mgrid_pos.z-rcoord.z); @@ -168,8 +168,8 @@ __global__ void set_phi_dphi_kernel( double dtmp = 0; const int it_nw = atom_type * nwmax; int iw_nr = it_nw * nrmax + ip; - int phi_idx = atoms_phi_start[atom_id + pre_atoms_num] + - bgrids_phi_len[bgrid_id] * mgrid_id; + int phi_idx = atom_phi_start[atom_id + pre_atoms_num] + + bgrid_phi_len[bgrid_id] * mgrid_id; for (int iw = 0; iw < atom_nw[atom_type]; iw++, iw_nr += nrmax) { if (atom_iw2_new[it_nw + iw]) @@ -199,8 +199,8 @@ __global__ void set_phi_dphi_kernel( } else { - int phi_idx = atoms_phi_start[atom_id + pre_atoms_num] + - bgrids_phi_len[bgrid_id] * mgrid_id; + int phi_idx = atom_phi_start[atom_id + pre_atoms_num] + + bgrid_phi_len[bgrid_id] * mgrid_id; for (int iw = 0; iw < atom_nw[atom_type]; iw++) { if(phi != nullptr) @@ -233,10 +233,10 @@ __global__ void set_ddphi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, double* __restrict__ ddphi_xx, double* __restrict__ ddphi_xy, double* __restrict__ ddphi_xz, @@ -253,15 +253,15 @@ __global__ void set_ddphi_kernel( for (int atom_id = threadIdx.x; atom_id < atoms_num; atom_id += blockDim.x) { const int atom_type = iat2it[atoms_iat[atom_id + pre_atoms_num]]; - const double3 rcoord = atoms_bgrids_rcoords[atom_id + pre_atoms_num]; + const double3 rcoord = atom_rcoords[atom_id + pre_atoms_num]; double coord[3]{mgrid_pos.x-rcoord.x, mgrid_pos.y-rcoord.y, mgrid_pos.z-rcoord.z}; double dist = norm3d(coord[0], coord[1], coord[2]); if (dist < rcut[atom_type]) { - int phi_idx = atoms_phi_start[atom_id + pre_atoms_num] + - bgrids_phi_len[bgrid_id] * mgrid_id; + int phi_idx = atom_phi_start[atom_id + pre_atoms_num] + + bgrid_phi_len[bgrid_id] * mgrid_id; for(int i = 0; i < 6; i++) { coord[i/2] += std::pow(-1, i%2) * 0.0001; @@ -362,17 +362,17 @@ __global__ void phi_mul_vldr3_kernel( const Real dr3, const Real* __restrict__ phi, const int mgrids_per_bgrid, - const int* __restrict__ mgrids_local_idx, - const int* __restrict__ bgrids_phi_len, - const int* __restrict__ bgrids_phi_start, + const int* __restrict__ mgrid_lidx, + const int* __restrict__ bgrid_phi_len, + const int* __restrict__ bgrid_phi_start, Real* __restrict__ result) { const int bgrid_id = blockIdx.y; const int mgrid_id = blockIdx.x; - const int phi_len = bgrids_phi_len[bgrid_id]; - const int phi_start = bgrids_phi_start[bgrid_id] + mgrid_id * phi_len; - const int mgrid_id_in_batch = bgrid_id * mgrids_per_bgrid + mgrid_id; - const Real vldr3 = vl[mgrids_local_idx[mgrid_id_in_batch]] * dr3; + const int phi_len = bgrid_phi_len[bgrid_id]; + const int phi_start = bgrid_phi_start[bgrid_id] + mgrid_id * phi_len; + const int batch_mgrid_id = bgrid_id * mgrids_per_bgrid + mgrid_id; + const Real vldr3 = vl[mgrid_lidx[batch_mgrid_id]] * dr3; for(int i = threadIdx.x; i < phi_len; i += blockDim.x) { result[phi_start + i] = phi[phi_start + i] * vldr3; @@ -394,20 +394,20 @@ __global__ void phi_dot_phi_kernel( const Real* __restrict__ phi_i, const Real* __restrict__ phi_j, const int mgrids_per_bgrid, - const int* __restrict__ mgrids_local_idx, - const int* __restrict__ bgrids_phi_len, - const int* __restrict__ bgrids_phi_start, + const int* __restrict__ mgrid_lidx, + const int* __restrict__ bgrid_phi_len, + const int* __restrict__ bgrid_phi_start, Real* __restrict__ rho) { __shared__ Real s_data[32]; // the length of s_data equals the max warp num of a block const int bgrid_id = blockIdx.y; const int mgrid_id = blockIdx.x; - const int phi_len = bgrids_phi_len[bgrid_id]; - const int phi_start = bgrids_phi_start[bgrid_id] + mgrid_id * phi_len; + const int phi_len = bgrid_phi_len[bgrid_id]; + const int phi_start = bgrid_phi_start[bgrid_id] + mgrid_id * phi_len; const Real* phi_i_mgrid = phi_i + phi_start; const Real* phi_j_mgrid = phi_j + phi_start; - const int mgrid_id_in_batch = bgrid_id * mgrids_per_bgrid + mgrid_id; - const int mgrid_local_idx = mgrids_local_idx[mgrid_id_in_batch]; + const int batch_mgrid_id = bgrid_id * mgrids_per_bgrid + mgrid_id; + const int mgrid_local_idx = mgrid_lidx[batch_mgrid_id]; const int tid = threadIdx.x; const int warp_id = tid / 32; const int lane_id = tid % 32; @@ -452,9 +452,9 @@ __global__ void phi_dot_dphi_kernel( const double* __restrict__ dphi_y, const double* __restrict__ dphi_z, const int mgrids_per_bgrid, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ bgrid_phi_len, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, + const int* __restrict__ atom_phi_start, const int* __restrict__ atoms_iat, const int* __restrict__ iat2it, const int* __restrict__ atom_nw, @@ -464,20 +464,20 @@ __global__ void phi_dot_dphi_kernel( const int bgrid_id = blockIdx.y; const int atoms_num = atoms_num_info[bgrid_id].x; const int pre_atoms_num = atoms_num_info[bgrid_id].y; - const int bgrid_phi_len = bgrids_phi_len[bgrid_id]; + const int b_phi_len = bgrid_phi_len[bgrid_id]; const int tid = threadIdx.x; const int warp_id = tid / 32; const int lane_id = tid % 32; for (int atom_id = blockIdx.x; atom_id < atoms_num; atom_id += gridDim.x) { - const int atom_phi_start = atoms_phi_start[atom_id + pre_atoms_num]; + const int a_phi_start = atom_phi_start[atom_id + pre_atoms_num]; const int iat = atoms_iat[atom_id + pre_atoms_num]; const int nw = atom_nw[iat2it[iat]]; double f[3] = {0.0, 0.0, 0.0}; for (int mgrid_id = 0; mgrid_id < mgrids_per_bgrid; mgrid_id++) { - const int phi_start = atom_phi_start + mgrid_id * bgrid_phi_len; + const int phi_start = a_phi_start + mgrid_id * b_phi_len; for (int iw = tid; iw < nw; iw += blockDim.x) { int phi_idx = phi_start + iw; @@ -529,11 +529,11 @@ __global__ void phi_dot_dphi_r_kernel( const double* __restrict__ dphi_y, const double* __restrict__ dphi_z, const int mgrids_per_bgrid, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ bgrid_phi_len, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, + const int* __restrict__ atom_phi_start, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const double3* __restrict__ mgrids_pos, const int* __restrict__ iat2it, const int* __restrict__ atom_nw, @@ -544,26 +544,26 @@ __global__ void phi_dot_dphi_r_kernel( const int bgrid_id = blockIdx.y; const int atoms_num = atoms_num_info[bgrid_id].x; const int pre_atoms_num = atoms_num_info[bgrid_id].y; - const int bgrid_phi_len = bgrids_phi_len[bgrid_id]; + const int b_phi_len = bgrid_phi_len[bgrid_id]; const int warp_id = tid / 32; const int lane_id = tid % 32; - + double stress[6]{0.0}; for (int mgrid_id = blockIdx.x; mgrid_id < mgrids_per_bgrid; mgrid_id += gridDim.x) { const double3 mgrid_pos = mgrids_pos[mgrid_id]; for (int atom_id = 0; atom_id < atoms_num; atom_id++) { - const int atom_phi_start = atoms_phi_start[atom_id + pre_atoms_num] + mgrid_id * bgrid_phi_len; + const int phi_start = atom_phi_start[atom_id + pre_atoms_num] + mgrid_id * b_phi_len; const int iat = atoms_iat[atom_id + pre_atoms_num]; const int nw = atom_nw[iat2it[iat]]; - const double3 rcoord = atoms_bgrids_rcoords[atom_id + pre_atoms_num]; // rcoord is the ralative coordinate of an atom and a biggrid + const double3 rcoord = atom_rcoords[atom_id + pre_atoms_num]; // rcoord is the ralative coordinate of an atom and a biggrid const double3 coord = make_double3(mgrid_pos.x-rcoord.x, // coord is the relative coordinate of an atom and a meshgrid mgrid_pos.y-rcoord.y, mgrid_pos.z-rcoord.z); for (int iw = tid; iw < nw; iw += blockDim.x) { - int phi_idx = atom_phi_start + iw; + int phi_idx = phi_start + iw; stress[0] += phi[phi_idx] * dphi_x[phi_idx] * coord.x; stress[1] += phi[phi_idx] * dphi_x[phi_idx] * coord.y; stress[2] += phi[phi_idx] * dphi_x[phi_idx] * coord.z; diff --git a/source/source_lcao/module_gint/kernel/phi_operator_kernel.cuh b/source/source_lcao/module_gint/kernel/phi_operator_kernel.cuh index 134216d45b8..9e51ae97c19 100644 --- a/source/source_lcao/module_gint/kernel/phi_operator_kernel.cuh +++ b/source/source_lcao/module_gint/kernel/phi_operator_kernel.cuh @@ -22,10 +22,10 @@ __global__ void set_phi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, Real* __restrict__ phi); __global__ void set_phi_dphi_kernel( @@ -44,10 +44,10 @@ __global__ void set_phi_dphi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, double* __restrict__ phi, double* __restrict__ dphi_x, double* __restrict__ dphi_y, @@ -69,10 +69,10 @@ __global__ void set_ddphi_kernel( const double* __restrict__ dpsi_u, const double3* __restrict__ mgrids_pos, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ atom_phi_start, + const int* __restrict__ bgrid_phi_len, double* __restrict__ ddphi_xx, double* __restrict__ ddphi_xy, double* __restrict__ ddphi_xz, @@ -86,9 +86,9 @@ __global__ void phi_mul_vldr3_kernel( const Real dr3, const Real* __restrict__ phi, const int mgrids_per_bgrid, - const int* __restrict__ mgrids_local_idx, - const int* __restrict__ bgrids_phi_len, - const int* __restrict__ bgrids_phi_start, + const int* __restrict__ mgrid_lidx, + const int* __restrict__ bgrid_phi_len, + const int* __restrict__ bgrid_phi_start, Real* __restrict__ result); // rho(ir) = \sum_{iwt} \phi_i(ir,iwt) * \phi_j^*(ir,iwt) @@ -98,9 +98,9 @@ __global__ void phi_dot_phi_kernel( const Real* __restrict__ phi_i, // phi_i(ir,iwt) const Real* __restrict__ phi_j, // phi_j(ir,iwt) const int mgrids_per_bgrid, // the number of mgrids of each biggrid - const int* __restrict__ mgrids_local_idx, // the idx of mgrid in local cell - const int* __restrict__ bgrids_phi_len, // the length of phi on a mgrid of a biggrid - const int* __restrict__ bgrids_phi_start, // the start idx in phi of each biggrid + const int* __restrict__ mgrid_lidx, // the idx of mgrid in local cell + const int* __restrict__ bgrid_phi_len, // the length of phi on a mgrid of a biggrid + const int* __restrict__ bgrid_phi_start, // the start idx in phi of each biggrid Real* __restrict__ rho); // rho(ir) __global__ void phi_dot_dphi_kernel( @@ -109,9 +109,9 @@ __global__ void phi_dot_dphi_kernel( const double* __restrict__ dphi_y, const double* __restrict__ dphi_z, const int mgrids_per_bgrid, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ bgrid_phi_len, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, + const int* __restrict__ atom_phi_start, const int* __restrict__ atoms_iat, const int* __restrict__ iat2it, const int* __restrict__ atom_nw, @@ -123,11 +123,11 @@ __global__ void phi_dot_dphi_r_kernel( const double* __restrict__ dphi_y, const double* __restrict__ dphi_z, const int mgrids_per_bgrid, - const int* __restrict__ bgrids_phi_len, + const int* __restrict__ bgrid_phi_len, const int2* __restrict__ atoms_num_info, - const int* __restrict__ atoms_phi_start, + const int* __restrict__ atom_phi_start, const int* __restrict__ atoms_iat, - const double3* __restrict__ atoms_bgrids_rcoords, + const double3* __restrict__ atom_rcoords, const double3* __restrict__ mgrids_pos, const int* __restrict__ iat2it, const int* __restrict__ atom_nw, diff --git a/source/source_lcao/module_gint/meshgrid_info.h b/source/source_lcao/module_gint/meshgrid_info.h index 6f91278756b..c1c276bc6e1 100644 --- a/source/source_lcao/module_gint/meshgrid_info.h +++ b/source/source_lcao/module_gint/meshgrid_info.h @@ -17,27 +17,27 @@ class MeshGridInfo meshgrid_vec2_(meshgrid_vec2), meshgrid_vec3_(meshgrid_vec3) { - // initialize the meshgrid_latvec0_ - meshgrid_latvec0_.e11 = meshgrid_vec1_.x; - meshgrid_latvec0_.e12 = meshgrid_vec1_.y; - meshgrid_latvec0_.e13 = meshgrid_vec1_.z; + // initialize the mgrid_latvec0_ + mgrid_latvec0_.e11 = meshgrid_vec1_.x; + mgrid_latvec0_.e12 = meshgrid_vec1_.y; + mgrid_latvec0_.e13 = meshgrid_vec1_.z; - meshgrid_latvec0_.e21 = meshgrid_vec2_.x; - meshgrid_latvec0_.e22 = meshgrid_vec2_.y; - meshgrid_latvec0_.e23 = meshgrid_vec2_.z; + mgrid_latvec0_.e21 = meshgrid_vec2_.x; + mgrid_latvec0_.e22 = meshgrid_vec2_.y; + mgrid_latvec0_.e23 = meshgrid_vec2_.z; - meshgrid_latvec0_.e31 = meshgrid_vec3_.x; - meshgrid_latvec0_.e32 = meshgrid_vec3_.y; - meshgrid_latvec0_.e33 = meshgrid_vec3_.z; + mgrid_latvec0_.e31 = meshgrid_vec3_.x; + mgrid_latvec0_.e32 = meshgrid_vec3_.y; + mgrid_latvec0_.e33 = meshgrid_vec3_.z; // initialize the GT matrix - meshgrid_GT_ = meshgrid_latvec0_.Inverse(); + meshgrid_GT_ = mgrid_latvec0_.Inverse(); - meshgrid_volume_ = std::abs(meshgrid_latvec0_.Det()); + mgrid_vol_ = std::abs(mgrid_latvec0_.Det()); } - double get_volume() const { return meshgrid_volume_; } - Vec3d get_cartesian_coord(const Vec3i& index_3d) const { return index_3d * meshgrid_latvec0_; } + double get_volume() const { return mgrid_vol_; } + Vec3d get_cartesian_coord(const Vec3i& index_3d) const { return index_3d * mgrid_latvec0_; } Vec3d get_direct_coord(const Vec3d& cart_coord) const { return cart_coord * meshgrid_GT_; } private: @@ -48,16 +48,16 @@ class MeshGridInfo // used to convert the (i, j, k) index of the meshgrid to the Cartesian coordinate // if meshrid_vec1_ is row vector, - // then meshgrid_latvec0_ = [meshgrid_vec1_; meshgrid_vec2_; meshgrid_vec3_], - // (i, j, k) * meshgrid_latvec0_ = (x, y, z) - Matrix3 meshgrid_latvec0_; + // then mgrid_latvec0_ = [meshgrid_vec1_; meshgrid_vec2_; meshgrid_vec3_], + // (i, j, k) * mgrid_latvec0_ = (x, y, z) + Matrix3 mgrid_latvec0_; // used to convert the Cartesian coordinate to the (i, j, k) index of the mesh grid - // meshgrid_GT_ = meshgrid_latvec0_.Inverse() + // meshgrid_GT_ = mgrid_latvec0_.Inverse() // (x, y, z) * meshgrid_GT_ = (i, j, k) Matrix3 meshgrid_GT_; - double meshgrid_volume_; + double mgrid_vol_; }; } // namespace ModuleGint \ No newline at end of file diff --git a/source/source_lcao/module_gint/phi_operator.cpp b/source/source_lcao/module_gint/phi_operator.cpp index 76dc5f8b4b2..7f5fa700a42 100644 --- a/source/source_lcao/module_gint/phi_operator.cpp +++ b/source/source_lcao/module_gint/phi_operator.cpp @@ -13,22 +13,22 @@ void PhiOperator::set_bgrid(std::shared_ptr biggrid) biggrid_->set_atoms_startidx(atoms_startidx_); biggrid_->set_atoms_phi_len(atoms_phi_len_); - biggrid_->set_mgrids_local_idx(meshgrids_local_idx_); + biggrid_->set_mgrids_local_idx(mgrid_lidx_); - // init is_atom_on_mgrid_ and atoms_relative_coords_ + // init is_atom_on_mgrid_ and atom_rcoords_ const int atoms_num = biggrid_->get_atoms_num(); - atoms_relative_coords_.resize(atoms_num); + atom_rcoords_.resize(atoms_num); is_atom_on_mgrid_.resize(biggrid_->get_mgrids_num() * atoms_num); for(int i = 0; i < atoms_num; ++i) { - biggrid_->set_atom_relative_coords(biggrid_->get_atom(i), atoms_relative_coords_[i]); + biggrid_->set_atom_relative_coords(biggrid_->get_atom(i), atom_rcoords_[i]); for(int j = 0; j < rows_; ++j) { - is_atom_on_mgrid_[i * rows_ + j] = atoms_relative_coords_[i][j].norm() <= biggrid_->get_atom(i)->get_rcut(); + is_atom_on_mgrid_[i * rows_ + j] = atom_rcoords_[i][j].norm() <= biggrid_->get_atom(i)->get_rcut(); } } - // init atom_pair_start_end_idx_ + // init atom_pair_range_ init_atom_pair_idx_(); } @@ -37,7 +37,7 @@ void PhiOperator::set_phi_dphi(double* phi, double* dphi_x, double* dphi_y, doub for(int i = 0; i < biggrid_->get_atoms_num(); ++i) { const auto atom = biggrid_->get_atom(i); - atom->set_phi_dphi(atoms_relative_coords_[i], cols_, phi, dphi_x, dphi_y, dphi_z); + atom->set_phi_dphi(atom_rcoords_[i], cols_, phi, dphi_x, dphi_y, dphi_z); if(phi != nullptr) { phi += atom->get_nw(); @@ -55,7 +55,7 @@ void PhiOperator::set_ddphi( for(int i = 0; i < biggrid_->get_atoms_num(); ++i) { const auto atom = biggrid_->get_atom(i); - atom->set_ddphi(atoms_relative_coords_[i], cols_, ddphi_xx, ddphi_xy, ddphi_xz, ddphi_yy, ddphi_yz, ddphi_zz); + atom->set_ddphi(atom_rcoords_[i], cols_, ddphi_xx, ddphi_xy, ddphi_xz, ddphi_yy, ddphi_yz, ddphi_zz); ddphi_xx += atom->get_nw(); ddphi_xy += atom->get_nw(); ddphi_xz += atom->get_nw(); @@ -108,7 +108,7 @@ void PhiOperator::phi_dot_dphi_r( for(int j = 0; j < biggrid_->get_atoms_num(); ++j) { const int start_idx = atoms_startidx_[j]; - const Vec3d& r3 = atoms_relative_coords_[j][i]; + const Vec3d& r3 = atom_rcoords_[j][i]; for(int k = 0; k < atoms_phi_len_[j]; ++k) { const int idx = i * cols_ + start_idx + k; @@ -151,7 +151,7 @@ void PhiOperator::cal_env_gamma( { tmp += phi[j * cols_ + start_idx + iw] * wfc[iw_lo]; } - rho[meshgrids_local_idx_[j]] += tmp; + rho[mgrid_lidx_[j]] += tmp; } } } @@ -204,7 +204,7 @@ void PhiOperator::cal_env_k( tmp += std::complex(phi[phi_start_idx + iw], 0.0) * wfc[iw_lo] * kphase; } } - rho[meshgrids_local_idx_[j]] += tmp.real(); + rho[mgrid_lidx_[j]] += tmp.real(); } } } @@ -217,7 +217,7 @@ void PhiOperator::cal_env_k( void PhiOperator::init_atom_pair_idx_() { int atoms_num = biggrid_->get_atoms_num(); - atom_pair_start_end_idx_.resize(atoms_num * (atoms_num + 1) / 2); + atom_pair_range_.resize(atoms_num * (atoms_num + 1) / 2); int mgrids_num = biggrid_->get_mgrids_num(); int atom_pair_idx = 0; for(int i = 0; i < atoms_num; ++i) @@ -243,8 +243,8 @@ void PhiOperator::init_atom_pair_idx_() break; } } - atom_pair_start_end_idx_[atom_pair_idx].first = start_idx; - atom_pair_start_end_idx_[atom_pair_idx].second = end_idx; + atom_pair_range_[atom_pair_idx].first = start_idx; + atom_pair_range_[atom_pair_idx].second = end_idx; atom_pair_idx++; } } diff --git a/source/source_lcao/module_gint/phi_operator.h b/source/source_lcao/module_gint/phi_operator.h index 323f99e8e93..83df19e9d26 100644 --- a/source/source_lcao/module_gint/phi_operator.h +++ b/source/source_lcao/module_gint/phi_operator.h @@ -19,7 +19,7 @@ namespace ModuleGint class PhiOperator { public: - enum class Triangular_Matrix{Upper, Lower, Full}; + enum class TriPart{Upper, Lower, Full}; // constructor PhiOperator()=default; @@ -70,7 +70,7 @@ class PhiOperator const T*const phi_i, // phi_i(ir,iwt) const T*const phi_j, // phi_j(ir,iwt) HContainer& hr, // hr(iwt_i,iwt_j) - const Triangular_Matrix triangular_matrix) const; + const TriPart part) const; // rho(ir) = \sum_{iwt} \phi_i(ir,iwt) * \phi_j(ir,iwt) template @@ -115,12 +115,12 @@ class PhiOperator void init_atom_pair_idx_(); // get the index of the first and the last meshgrid that both atom a and atom b affect - // Note that atom_pair_start_end_idx_ only stores the cases where a <= b, so this function is needed to retrieve the value + // Note that atom_pair_range_ only stores the cases where a <= b, so this function is needed to retrieve the value const std::pair& get_atom_pair_start_end_idx_(int a, int b) const { int x = std::min(a, b); int y = std::abs(a - b); - return atom_pair_start_end_idx_[(2 * biggrid_->get_atoms_num() - x + 1) * x / 2 + y]; + return atom_pair_range_[(2 * biggrid_->get_atoms_num() - x + 1) * x / 2 + y]; } bool is_atom_on_mgrid(int atom_idx, int mgrid_idx) const @@ -137,14 +137,14 @@ class PhiOperator int cols_; // the local index of the meshgrids - std::vector meshgrids_local_idx_; + std::vector mgrid_lidx_; // the big grid that the phi matrix is associated with std::shared_ptr biggrid_; // the relative coordinates of the atoms and the meshgrids - // atoms_relative_coords_[i][j] is the relative coordinate of the jth meshgrid and the ith atom - std::vector> atoms_relative_coords_; + // atom_rcoords_[i][j] is the relative coordinate of the jth meshgrid and the ith atom + std::vector> atom_rcoords_; // record whether the atom affects the meshgrid // is_atom_on_mgrid_[i * rows_ + j] = true if the ith atom affects jhe ith meshgrid, otherwise false @@ -159,7 +159,7 @@ class PhiOperator std::vector atoms_phi_len_; // This data structure is used to store the index of the first and last meshgrid affected by each atom pair - std::vector> atom_pair_start_end_idx_; + std::vector> atom_pair_range_; }; } diff --git a/source/source_lcao/module_gint/phi_operator.hpp b/source/source_lcao/module_gint/phi_operator.hpp index ee8696707e6..fefafb3f8f3 100644 --- a/source/source_lcao/module_gint/phi_operator.hpp +++ b/source/source_lcao/module_gint/phi_operator.hpp @@ -13,7 +13,7 @@ void PhiOperator::set_phi(T* phi) const for(int i = 0; i < biggrid_->get_atoms_num(); ++i) { const auto atom = biggrid_->get_atom(i); - atom->set_phi(atoms_relative_coords_[i], cols_, phi); + atom->set_phi(atom_rcoords_[i], cols_, phi); phi += atom->get_nw(); } } @@ -94,7 +94,7 @@ void PhiOperator::phi_mul_vldr3( int idx = 0; for(int i = 0; i < biggrid_->get_mgrids_num(); i++) { - T vldr3_mgrid = vl[meshgrids_local_idx_[i]] * dr3; + T vldr3_mgrid = vl[mgrid_lidx_[i]] * dr3; for(int j = 0; j < cols_; j++) { result[idx] = phi[idx] * vldr3_mgrid; @@ -110,7 +110,7 @@ void PhiOperator::phi_mul_phi( const T*const phi_i, // phi_i(ir,iwt) const T*const phi_j, // phi_j(ir,iwt) HContainer& hr, // hr(iwt_i,iwt_j) - const Triangular_Matrix triangular_matrix) const + const TriPart part) const { std::vector tmp_hr; for(int i = 0; i < biggrid_->get_atoms_num(); ++i) @@ -128,12 +128,12 @@ void PhiOperator::phi_mul_phi( const int n_j = atoms_phi_len_[j]; // only calculate the upper triangle matrix - if(triangular_matrix==Triangular_Matrix::Upper && iat_i>iat_j) + if(part==TriPart::Upper && iat_i>iat_j) { continue; } // only calculate the upper triangle matrix - else if(triangular_matrix==Triangular_Matrix::Lower && iat_iget_mgrids_num(); ++i) { - rho[meshgrids_local_idx_[i]] += static_cast( + rho[mgrid_lidx_[i]] += static_cast( BlasConnector::dotc(cols_, phi_j + i * cols_, inc, phi_i + i * cols_, inc)); } }