Remove all ctx parameters in set_memory_op

Critsium-xy · Critsium-xy · commit d485c505b387 · 2025-01-16T15:30:42.000+08:00
diff --git a/source/module_base/math_chebyshev.cpp b/source/module_base/math_chebyshev.cpp
@@ -437,7 +437,7 @@ void Chebyshev<REAL, Device>::calfinalvec_real(
     funA(arrayn_1, arrayn, m);
 
     // 0- & 1-st order
-    setmem_complex_op()(this->ctx, waveout, 0, ndmxt);
+    setmem_complex_op()(waveout, 0, ndmxt);
     std::complex<REAL> coef0 = std::complex<REAL>(coefr_cpu[0], 0);
     container::kernels::blas_axpy<std::complex<REAL>, ct_Device>()(ndmxt, &coef0, arrayn_1, 1, waveout, 1);
     std::complex<REAL> coef1 = std::complex<REAL>(coefr_cpu[1], 0);
@@ -505,7 +505,7 @@ void Chebyshev<REAL, Device>::calfinalvec_complex(
     funA(arrayn_1, arrayn, m);
 
     // 0- & 1-st order
-    setmem_complex_op()(this->ctx, waveout, 0, ndmxt);
+    setmem_complex_op()(waveout, 0, ndmxt);
     container::kernels::blas_axpy<std::complex<REAL>, ct_Device>()(ndmxt, &coefc_cpu[0], arrayn_1, 1, waveout, 1);
     container::kernels::blas_axpy<std::complex<REAL>, ct_Device>()(ndmxt, &coefc_cpu[1], arrayn, 1, waveout, 1);
     // for (int i = 0; i < ndmxt; ++i)
diff --git a/source/module_base/module_device/cuda/memory_op.cu b/source/module_base/module_device/cuda/memory_op.cu
@@ -53,7 +53,7 @@ __global__ void cast_memory(std::complex<FPTYPE_out>* out, const FPTYPE_in* in,
 
 template <typename FPTYPE>
 void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,
-                                                                     const size_t size,
+                                                                   const size_t size,
                                                                    const char* record_in)
 {
     if (arr != nullptr)
@@ -78,8 +78,7 @@ void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,
 }
 
 template <typename FPTYPE>
-void set_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* dev,
-                                                                FPTYPE* arr,
+void set_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE* arr,
                                                                 const int var,
                                                                 const size_t size)
 {
diff --git a/source/module_base/module_device/memory_op.cpp b/source/module_base/module_device/memory_op.cpp
@@ -45,7 +45,7 @@ struct resize_memory_op<FPTYPE, base_device::DEVICE_CPU>
 template <typename FPTYPE>
 struct set_memory_op<FPTYPE, base_device::DEVICE_CPU>
 {
-    void operator()(const base_device::DEVICE_CPU* dev, FPTYPE* arr, const int var, const size_t size)
+    void operator()(FPTYPE* arr, const int var, const size_t size)
     {
         ModuleBase::OMP_PARALLEL([&](int num_thread, int thread_id) {
             int beg = 0, len = 0;
@@ -166,7 +166,7 @@ struct resize_memory_op<FPTYPE, base_device::DEVICE_GPU>
 template <typename FPTYPE>
 struct set_memory_op<FPTYPE, base_device::DEVICE_GPU>
 {
-    void operator()(const base_device::DEVICE_GPU* dev, FPTYPE* arr, const int var, const size_t size)
+    void operator()(FPTYPE* arr, const int var, const size_t size)
     {
     }
 };
diff --git a/source/module_base/module_device/memory_op.h b/source/module_base/module_device/memory_op.h
@@ -32,13 +32,12 @@ struct set_memory_op
     /// @brief memset for multi-device
     ///
     /// Input Parameters
-    /// \param dev : the type of computing device
     /// \param var : the specified constant value
     /// \param size : array size
     ///
     /// Output Parameters
     /// \param arr : output array initialized by the input value
-    void operator()(const Device* dev, FPTYPE* arr, const int var, const size_t size);
+    void operator()(FPTYPE* arr, const int var, const size_t size);
 };
 
 template <typename FPTYPE, typename Device_out, typename Device_in>
@@ -120,7 +119,7 @@ struct resize_memory_op<FPTYPE, base_device::DEVICE_GPU>
 template <typename FPTYPE>
 struct set_memory_op<FPTYPE, base_device::DEVICE_GPU>
 {
-    void operator()(const base_device::DEVICE_GPU* dev, FPTYPE* arr, const int var, const size_t size);
+    void operator()(FPTYPE* arr, const int var, const size_t size);
 };
 
 template <typename FPTYPE>
diff --git a/source/module_base/module_device/rocm/memory_op.hip.cu b/source/module_base/module_device/rocm/memory_op.hip.cu
@@ -39,8 +39,7 @@ __global__ void cast_memory(std::complex<FPTYPE_out>* out, const std::complex<FP
 }
 
 template <typename FPTYPE>
-void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* dev,
-                                                                   FPTYPE*& arr,
+void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,
                                                                    const size_t size,
                                                                    const char* record_in)
 {
diff --git a/source/module_base/module_device/test/memory_test.cpp b/source/module_base/module_device/test/memory_test.cpp
@@ -91,7 +91,7 @@ class TestModulePsiMemory : public ::testing::Test
 TEST_F(TestModulePsiMemory, set_memory_op_double_cpu)
 {
     std::vector<double> v_xx = xx;
-    set_memory_double_cpu_op()(cpu_ctx, v_xx.data(), 0, xx.size());
+    set_memory_double_cpu_op()(v_xx.data(), 0, xx.size());
     for (int ii = 0; ii < xx.size(); ii++)
     {
         EXPECT_EQ(v_xx[ii], 0.0);
@@ -101,7 +101,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_double_cpu)
 TEST_F(TestModulePsiMemory, set_memory_op_complex_double_cpu)
 {
     std::vector<std::complex<double>> vz_xx = z_xx;
-    set_memory_complex_double_cpu_op()(cpu_ctx, vz_xx.data(), 0, z_xx.size());
+    set_memory_complex_double_cpu_op()(vz_xx.data(), 0, z_xx.size());
     for (int ii = 0; ii < z_xx.size(); ii++)
     {
         EXPECT_EQ(vz_xx[ii], std::complex<double>(0.0, 0.0));
@@ -175,7 +175,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_double_gpu)
 {
     thrust::device_ptr<double> d_xx = thrust::device_malloc<double>(xx.size());
     thrust::copy(xx.begin(), xx.end(), d_xx);
-    set_memory_double_gpu_op()(gpu_ctx, thrust::raw_pointer_cast(d_xx), 0, xx.size());
+    set_memory_double_gpu_op()(thrust::raw_pointer_cast(d_xx), 0, xx.size());
     thrust::host_vector<double> h_xx(xx.size());
     thrust::copy(d_xx, d_xx + xx.size(), h_xx.begin());
     for (int ii = 0; ii < xx.size(); ii++)
@@ -188,7 +188,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_complex_double_gpu)
 {
     thrust::device_ptr<std::complex<double>> dz_xx = thrust::device_malloc<std::complex<double>>(z_xx.size());
     thrust::copy(z_xx.begin(), z_xx.end(), dz_xx);
-    set_memory_complex_double_gpu_op()(gpu_ctx, thrust::raw_pointer_cast(dz_xx), 0, z_xx.size());
+    set_memory_complex_double_gpu_op()(thrust::raw_pointer_cast(dz_xx), 0, z_xx.size());
     thrust::host_vector<std::complex<double>> h_xx(z_xx.size());
     thrust::copy(dz_xx, dz_xx + z_xx.size(), h_xx.begin());
     for (int ii = 0; ii < z_xx.size(); ii++)
diff --git a/source/module_basis/module_pw/pw_transform_k.cpp b/source/module_basis/module_pw/pw_transform_k.cpp
@@ -413,7 +413,6 @@ void PW_Basis_K::recip_to_real(const base_device::DEVICE_GPU* ctx,
     assert(this->poolnproc == 1);
     // ModuleBase::GlobalFunc::ZEROS(fft_bundle.get_auxr_3d_data<float>(), this->nxyz);
     base_device::memory::set_memory_op<std::complex<float>, base_device::DEVICE_GPU>()(
-        ctx,
         this->fft_bundle.get_auxr_3d_data<float>(),
         0,
         this->nxyz);
@@ -450,7 +449,6 @@ void PW_Basis_K::recip_to_real(const base_device::DEVICE_GPU* ctx,
     assert(this->poolnproc == 1);
     // ModuleBase::GlobalFunc::ZEROS(fft_bundle.get_auxr_3d_data<double>(), this->nxyz);
     base_device::memory::set_memory_op<std::complex<double>, base_device::DEVICE_GPU>()(
-        ctx,
         this->fft_bundle.get_auxr_3d_data<double>(),
         0,
         this->nxyz);
diff --git a/source/module_elecstate/elecstate_pw.cpp b/source/module_elecstate/elecstate_pw.cpp
@@ -118,15 +118,15 @@ void ElecStatePW<T, Device>::psiToRho(const psi::Psi<T, Device>& psi)
 	{
         // denghui replaced at 20221110
 		// ModuleBase::GlobalFunc::ZEROS(this->rho[is], this->charge->nrxx);
-        setmem_var_op()(this->ctx, this->rho[is], 0,  this->charge->nrxx);
+        setmem_var_op()(this->rho[is], 0,  this->charge->nrxx);
         if (get_xc_func_type() == 3)
         {
             // ModuleBase::GlobalFunc::ZEROS(this->charge->kin_r[is], this->charge->nrxx);
-            setmem_var_op()(this->ctx, this->kin_r[is], 0,  this->charge->nrxx);
+            setmem_var_op()(this->kin_r[is], 0,  this->charge->nrxx);
         }
         if (PARAM.globalv.double_grid || PARAM.globalv.use_uspp)
         {
-            setmem_complex_op()(this->ctx, this->rhog[is], 0, this->charge->rhopw->npw);
+            setmem_complex_op()(this->rhog[is], 0, this->charge->rhopw->npw);
         }
     }
 
@@ -244,7 +244,7 @@ void ElecStatePW<T, Device>::rhoBandK(const psi::Psi<T, Device>& psi)
             {
                 for (int j = 0; j < 3; j++)
                 {
-                    setmem_complex_op()(this->ctx, this->wfcr, 0,  this->charge->nrxx);
+                    setmem_complex_op()(this->wfcr, 0,  this->charge->nrxx);
 
                     meta_op()(this->ctx,
                               ik,
@@ -280,7 +280,7 @@ void ElecStatePW<T, Device>::cal_becsum(const psi::Psi<T, Device>& psi)
     resmem_complex_op()(becp, nbands * nkb, "ElecState<PW>::becp");
     const int nh_tot = this->ppcell->nhm * (this->ppcell->nhm + 1) / 2;
     resmem_var_op()(becsum, nh_tot * ucell->nat * PARAM.inp.nspin, "ElecState<PW>::becsum");
-    setmem_var_op()(this->ctx, becsum, 0, nh_tot * ucell->nat * PARAM.inp.nspin);
+    setmem_var_op()(becsum, 0, nh_tot * ucell->nat * PARAM.inp.nspin);
 
     for (int ik = 0; ik < psi.get_nk(); ++ik)
     {
diff --git a/source/module_elecstate/elecstate_pw_cal_tau.cpp b/source/module_elecstate/elecstate_pw_cal_tau.cpp
@@ -9,7 +9,7 @@ void ElecStatePW<T, Device>::cal_tau(const psi::Psi<T, Device>& psi)
     ModuleBase::TITLE("ElecStatePW", "cal_tau");
     for(int is=0; is<PARAM.inp.nspin; is++)
 	{
-        setmem_var_op()(this->ctx, this->kin_r[is], 0,  this->charge->nrxx);
+        setmem_var_op()(this->kin_r[is], 0,  this->charge->nrxx);
 	}
 
     for (int ik = 0; ik < psi.get_nk(); ++ik)
@@ -31,7 +31,7 @@ void ElecStatePW<T, Device>::cal_tau(const psi::Psi<T, Device>& psi)
             // kinetic energy density
             for (int j = 0; j < 3; j++)
             {
-                setmem_complex_op()(this->ctx, this->wfcr, 0,  this->charge->nrxx);
+                setmem_complex_op()(this->wfcr, 0,  this->charge->nrxx);
 
                 meta_op()(this->ctx,
                             ik,
diff --git a/source/module_elecstate/elecstate_pw_sdft.cpp b/source/module_elecstate/elecstate_pw_sdft.cpp
@@ -16,7 +16,7 @@ void ElecStatePW_SDFT<T, Device>::psiToRho(const psi::Psi<T, Device>& psi)
     const int nspin = PARAM.inp.nspin;
     for (int is = 0; is < nspin; is++)
     {
-        setmem_var_op()(this->ctx, this->rho[is], 0, this->charge->nrxx);
+        setmem_var_op()(this->rho[is], 0, this->charge->nrxx);
     }
 
     if (GlobalV::MY_STOGROUP == 0)
diff --git a/source/module_hamilt_pw/hamilt_pwdft/VNL_in_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/VNL_in_pw.cpp
@@ -1080,7 +1080,7 @@ void pseudopot_cell_vnl::radial_fft_q(Device* ctx,
     const int ivl = nhtolm(itype, ih);
     const int jvl = nhtolm(itype, jh);
 
-    setmem_complex_op()(ctx, qg, 0, ng);
+    setmem_complex_op()(qg, 0, ng);
 
     const double* qnorm_double = reinterpret_cast<const double*>(qnorm);
 
diff --git a/source/module_hamilt_pw/hamilt_pwdft/forces_nl.cpp b/source/module_hamilt_pw/hamilt_pwdft/forces_nl.cpp
@@ -28,7 +28,7 @@ void Forces<FPTYPE, Device>::cal_force_nl(ModuleBase::matrix& forcenl,
     // allocate memory for the force
     FPTYPE* force = nullptr;
     resmem_var_op()(force, ucell_in.nat * 3);
-    base_device::memory::set_memory_op<FPTYPE, Device>()(this->ctx, force, 0.0, ucell_in.nat * 3);
+    base_device::memory::set_memory_op<FPTYPE, Device>()(force, 0.0, ucell_in.nat * 3);
 
     hamilt::FS_Nonlocal_tools<FPTYPE, Device> nl_tools(&nlpp, &ucell_in, p_kv, wfc_basis, p_sf, wg, &ekb);
 
diff --git a/source/module_hamilt_pw/hamilt_pwdft/forces_onsite.cpp b/source/module_hamilt_pw/hamilt_pwdft/forces_onsite.cpp
@@ -24,7 +24,7 @@ void Forces<FPTYPE, Device>::cal_force_onsite(ModuleBase::matrix& force_onsite,
     // allocate memory for the force
     FPTYPE* force = nullptr;
     resmem_var_op()(force, ucell_in.nat * 3);
-    base_device::memory::set_memory_op<FPTYPE, Device>()(this->ctx, force, 0.0, ucell_in.nat * 3);
+    base_device::memory::set_memory_op<FPTYPE, Device>()(force, 0.0, ucell_in.nat * 3);
 
     auto* onsite_p = projectors::OnsiteProjector<FPTYPE, Device>::get_instance();
 
diff --git a/source/module_hamilt_pw/hamilt_pwdft/hamilt_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/hamilt_pw.cpp
@@ -295,7 +295,7 @@ void HamiltPW<T, Device>::sPsi(const T* psi_in, // psi
         }
 
         resmem_complex_op()(ps, this->ppcell->nkb * nbands, "Hamilt<PW>::ps");
-        setmem_complex_op()(this->ctx, ps, 0, this->ppcell->nkb * nbands);
+        setmem_complex_op()(ps, 0, this->ppcell->nkb * nbands);
 
         // spsi = psi + sum qq <beta|psi> |beta>
         if (PARAM.inp.noncolin)
diff --git a/source/module_hamilt_pw/hamilt_pwdft/operator_pw/meta_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/operator_pw/meta_pw.cpp
@@ -55,7 +55,7 @@ void Meta<OperatorPW<T, Device>>::act(
     ModuleBase::timer::tick("Operator", "MetaPW");
     if(is_first_node)
     {
-        setmem_complex_op()(this->ctx, tmhpsi, 0, nbasis*nbands/npol);
+        setmem_complex_op()(tmhpsi, 0, nbasis*nbands/npol);
     }
 
     const int current_spin = this->isk[this->ik];
diff --git a/source/module_hamilt_pw/hamilt_pwdft/operator_pw/nonlocal_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/operator_pw/nonlocal_pw.cpp
@@ -75,7 +75,7 @@ void Nonlocal<OperatorPW<T, Device>>::add_nonlocal_pp(T *hpsi_in, const T *becp,
         resmem_complex_op()(this->ps, nkb * m, "Nonlocal<PW>::ps");
         this->nkb_m = m * nkb;
     }
-    setmem_complex_op()(this->ctx, this->ps, 0, nkb * m);
+    setmem_complex_op()(this->ps, 0, nkb * m);
 
     int sum = 0;
     int iat = 0;
@@ -221,7 +221,7 @@ void Nonlocal<OperatorPW<T, Device>>::act(
     ModuleBase::timer::tick("Operator", "NonlocalPW");
     if(is_first_node)
     {
-        setmem_complex_op()(this->ctx, tmhpsi, 0, nbasis*nbands/npol);
+        setmem_complex_op()(tmhpsi, 0, nbasis*nbands/npol);
     }
     if(!PARAM.inp.use_paw)
     {
diff --git a/source/module_hamilt_pw/hamilt_pwdft/operator_pw/onsite_proj_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/operator_pw/onsite_proj_pw.cpp
@@ -130,7 +130,7 @@ void OnsiteProj<OperatorPW<T, Device>>::cal_ps_delta_spin(const int npol, const
         resmem_complex_op()(this->ps, tnp * m, "OnsiteProj<PW>::ps");
         this->nkb_m = m * tnp;
     }
-    setmem_complex_op()(this->ctx, this->ps, 0, tnp * m);
+    setmem_complex_op()(this->ps, 0, tnp * m);
 
     if(!this->init_delta_spin)
     {
@@ -230,7 +230,7 @@ void OnsiteProj<OperatorPW<T, Device>>::cal_ps_dftu(const int npol, const int m)
     }
     if(!this->has_delta_spin) 
     {
-        setmem_complex_op()(this->ctx, this->ps, 0, tnp * m);
+        setmem_complex_op()(this->ps, 0, tnp * m);
     }
 
     if(!this->init_dftu)
diff --git a/source/module_hamilt_pw/hamilt_pwdft/operator_pw/veff_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/operator_pw/veff_pw.cpp
@@ -48,7 +48,7 @@ void Veff<OperatorPW<T, Device>>::act(
     ModuleBase::timer::tick("Operator", "VeffPW");
     if(is_first_node)
     {
-        setmem_complex_op()(this->ctx, tmhpsi, 0, nbasis*nbands/npol);
+        setmem_complex_op()(tmhpsi, 0, nbasis*nbands/npol);
     }
 
     int max_npw = nbasis / npol;
diff --git a/source/module_hamilt_pw/hamilt_pwdft/stress_func_nl.cpp b/source/module_hamilt_pw/hamilt_pwdft/stress_func_nl.cpp
@@ -31,7 +31,7 @@ void Stress_Func<FPTYPE, Device>::stress_nl(ModuleBase::matrix& sigma,
 
     FPTYPE* stress_device = nullptr;
     resmem_var_op()(stress_device, 9);
-    setmem_var_op()(this->ctx, stress_device, 0, 9);
+    setmem_var_op()(stress_device, 0, 9);
     std::vector<FPTYPE> sigmanlc(9, 0.0);
 
     hamilt::FS_Nonlocal_tools<FPTYPE, Device> nl_tools(&nlpp_in, &ucell_in, p_kv, wfc_basis, p_sf, wg, &ekb);
diff --git a/source/module_hamilt_pw/hamilt_pwdft/stress_func_onsite.cpp b/source/module_hamilt_pw/hamilt_pwdft/stress_func_onsite.cpp
@@ -23,7 +23,7 @@ void Stress_Func<FPTYPE, Device>::stress_onsite(ModuleBase::matrix& sigma,
 
     FPTYPE* stress_device = nullptr;
     resmem_var_op()(stress_device, 9);
-    setmem_var_op()(this->ctx, stress_device, 0, 9);
+    setmem_var_op()(stress_device, 0, 9);
     std::vector<FPTYPE> sigma_onsite(9, 0.0);
 
     auto* onsite_p = projectors::OnsiteProjector<FPTYPE, Device>::get_instance();
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_forces.cpp b/source/module_hamilt_pw/hamilt_stodft/sto_forces.cpp
@@ -218,7 +218,7 @@ void Sto_Forces<FPTYPE, Device>::cal_sto_force_nl(
    // allocate memory for the force
     FPTYPE* force = nullptr;
     resmem_var_op()(force, ucell.nat * 3);
-    base_device::memory::set_memory_op<FPTYPE, Device>()(this->ctx, force, 0.0, ucell.nat * 3);
+    base_device::memory::set_memory_op<FPTYPE, Device>()(force, 0.0, ucell.nat * 3);
 
     hamilt::FS_Nonlocal_tools<FPTYPE, Device> nl_tools(&nlpp, &ucell, p_kv, wfc_basis, p_sf, wg, nullptr);
 
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_iter.cpp b/source/module_hamilt_pw/hamilt_stodft/sto_iter.cpp
@@ -362,7 +362,7 @@ void Stochastic_Iter<T, Device>::calPn(const int& ik, Stochastic_WF<T, Device>&
         }
         else
         {
-            setmem_var_op()(this->ctx, spolyv, 0, norder * norder);
+            setmem_var_op()(spolyv, 0, norder * norder);
         }
     }
     T* pchi;
@@ -597,7 +597,7 @@ void Stochastic_Iter<T, Device>::cal_storho(const UnitCell& ucell,
     }
     for (int is = 0; is < nspin; is++)
     {
-        setmem_var_op()(this->ctx, pes->rho[is], 0, nrxx);
+        setmem_var_op()(pes->rho[is], 0, nrxx);
     }
     for (int ik = 0; ik < this->pkv->get_nks(); ++ik)
     {
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_stress_pw.cpp b/source/module_hamilt_pw/hamilt_stodft/sto_stress_pw.cpp
@@ -168,7 +168,7 @@ void Sto_Stress_PW<FPTYPE, Device>::sto_stress_nl(ModuleBase::matrix& sigma,
     // allocate memory for the stress
     FPTYPE* stress_device = nullptr;
     resmem_var_op()(stress_device, 9);
-    setmem_var_op()(this->ctx, stress_device, 0, 9);
+    setmem_var_op()(stress_device, 0, 9);
     std::vector<FPTYPE> sigmanlc(9, 0.0);
 
     hamilt::FS_Nonlocal_tools<FPTYPE, Device> nl_tools(&nlpp, &ucell, p_kv, wfc_basis, p_sf, wg, nullptr);
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_wf.cpp b/source/module_hamilt_pw/hamilt_stodft/sto_wf.cpp
@@ -52,7 +52,7 @@ void Stochastic_WF<T, Device>::allocate_chiallorder(const int& norder)
     for (int ik = 0; ik < this->nks; ++ik)
     {
         chiallorder[ik].resize(1, this->nchip[ik] * this->npwx, norder);
-        setmem_complex_op()(chiallorder[ik].get_device(), chiallorder[ik].get_pointer(), 0, chiallorder[ik].size());
+        setmem_complex_op()(chiallorder[ik].get_pointer(), 0, chiallorder[ik].size());
     }
 }
 
diff --git a/source/module_hsolver/diago_dav_subspace.cpp b/source/module_hsolver/diago_dav_subspace.cpp
diff --git a/source/module_hsolver/diago_david.cpp b/source/module_hsolver/diago_david.cpp
diff --git a/source/module_hsolver/diago_iter_assist.cpp b/source/module_hsolver/diago_iter_assist.cpp
diff --git a/source/module_psi/psi.cpp b/source/module_psi/psi.cpp

Original file line number	Diff line number	Diff line change
`@@ -53,7 +53,7 @@ __global__ void cast_memory(std::complex<FPTYPE_out>* out, const FPTYPE_in* in,`
`53`	`53`
`54`	`54`	`template <typename FPTYPE>`
`55`	`55`	`void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,`
`56`		`- const size_t size,`
	`56`	`+ const size_t size,`
`57`	`57`	`const char* record_in)`
`58`	`58`	`{`
`59`	`59`	`if (arr != nullptr)`
`@@ -78,8 +78,7 @@ void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,`
`78`	`78`	`}`
`79`	`79`
`80`	`80`	`template <typename FPTYPE>`
`81`		`-void set_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* dev,`
`82`		`- FPTYPE* arr,`
	`81`	`+void set_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE* arr,`
`83`	`82`	`const int var,`
`84`	`83`	`const size_t size)`
`85`	`84`	`{`
Original file line number	Diff line number	Diff line change
`@@ -45,7 +45,7 @@ struct resize_memory_op<FPTYPE, base_device::DEVICE_CPU>`
`45`	`45`	`template <typename FPTYPE>`
`46`	`46`	`struct set_memory_op<FPTYPE, base_device::DEVICE_CPU>`
`47`	`47`	`{`
`48`		`- void operator()(const base_device::DEVICE_CPU* dev, FPTYPE* arr, const int var, const size_t size)`
	`48`	`+ void operator()(FPTYPE* arr, const int var, const size_t size)`
`49`	`49`	`{`
`50`	`50`	`ModuleBase::OMP_PARALLEL([&](int num_thread, int thread_id) {`
`51`	`51`	`int beg = 0, len = 0;`
`@@ -166,7 +166,7 @@ struct resize_memory_op<FPTYPE, base_device::DEVICE_GPU>`
`166`	`166`	`template <typename FPTYPE>`
`167`	`167`	`struct set_memory_op<FPTYPE, base_device::DEVICE_GPU>`
`168`	`168`	`{`
`169`		`- void operator()(const base_device::DEVICE_GPU* dev, FPTYPE* arr, const int var, const size_t size)`
	`169`	`+ void operator()(FPTYPE* arr, const int var, const size_t size)`
`170`	`170`	`{`
`171`	`171`	`}`
`172`	`172`	`};`
Original file line number	Diff line number	Diff line change
`@@ -39,8 +39,7 @@ __global__ void cast_memory(std::complex<FPTYPE_out>* out, const std::complex<FP`
`39`	`39`	`}`
`40`	`40`
`41`	`41`	`template <typename FPTYPE>`
`42`		`-void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* dev,`
`43`		`- FPTYPE*& arr,`
	`42`	`+void resize_memory_op<FPTYPE, base_device::DEVICE_GPU>::operator()(FPTYPE*& arr,`
`44`	`43`	`const size_t size,`
`45`	`44`	`const char* record_in)`
`46`	`45`	`{`
Original file line number	Diff line number	Diff line change
`@@ -91,7 +91,7 @@ class TestModulePsiMemory : public ::testing::Test`
`91`	`91`	`TEST_F(TestModulePsiMemory, set_memory_op_double_cpu)`
`92`	`92`	`{`
`93`	`93`	`std::vector<double> v_xx = xx;`
`94`		`- set_memory_double_cpu_op()(cpu_ctx, v_xx.data(), 0, xx.size());`
	`94`	`+ set_memory_double_cpu_op()(v_xx.data(), 0, xx.size());`
`95`	`95`	`for (int ii = 0; ii < xx.size(); ii++)`
`96`	`96`	`{`
`97`	`97`	`EXPECT_EQ(v_xx[ii], 0.0);`
`@@ -101,7 +101,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_double_cpu)`
`101`	`101`	`TEST_F(TestModulePsiMemory, set_memory_op_complex_double_cpu)`
`102`	`102`	`{`
`103`	`103`	`std::vector<std::complex<double>> vz_xx = z_xx;`
`104`		`- set_memory_complex_double_cpu_op()(cpu_ctx, vz_xx.data(), 0, z_xx.size());`
	`104`	`+ set_memory_complex_double_cpu_op()(vz_xx.data(), 0, z_xx.size());`
`105`	`105`	`for (int ii = 0; ii < z_xx.size(); ii++)`
`106`	`106`	`{`
`107`	`107`	`EXPECT_EQ(vz_xx[ii], std::complex<double>(0.0, 0.0));`
`@@ -175,7 +175,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_double_gpu)`
`175`	`175`	`{`
`176`	`176`	`thrust::device_ptr<double> d_xx = thrust::device_malloc<double>(xx.size());`
`177`	`177`	`thrust::copy(xx.begin(), xx.end(), d_xx);`
`178`		`- set_memory_double_gpu_op()(gpu_ctx, thrust::raw_pointer_cast(d_xx), 0, xx.size());`
	`178`	`+ set_memory_double_gpu_op()(thrust::raw_pointer_cast(d_xx), 0, xx.size());`
`179`	`179`	`thrust::host_vector<double> h_xx(xx.size());`
`180`	`180`	`thrust::copy(d_xx, d_xx + xx.size(), h_xx.begin());`
`181`	`181`	`for (int ii = 0; ii < xx.size(); ii++)`
`@@ -188,7 +188,7 @@ TEST_F(TestModulePsiMemory, set_memory_op_complex_double_gpu)`
`188`	`188`	`{`
`189`	`189`	`thrust::device_ptr<std::complex<double>> dz_xx = thrust::device_malloc<std::complex<double>>(z_xx.size());`
`190`	`190`	`thrust::copy(z_xx.begin(), z_xx.end(), dz_xx);`
`191`		`- set_memory_complex_double_gpu_op()(gpu_ctx, thrust::raw_pointer_cast(dz_xx), 0, z_xx.size());`
	`191`	`+ set_memory_complex_double_gpu_op()(thrust::raw_pointer_cast(dz_xx), 0, z_xx.size());`
`192`	`192`	`thrust::host_vector<std::complex<double>> h_xx(z_xx.size());`
`193`	`193`	`thrust::copy(dz_xx, dz_xx + z_xx.size(), h_xx.begin());`
`194`	`194`	`for (int ii = 0; ii < z_xx.size(); ii++)`
Original file line number	Diff line number	Diff line change
`@@ -118,15 +118,15 @@ void ElecStatePW<T, Device>::psiToRho(const psi::Psi<T, Device>& psi)`
`118`	`118`	`{`
`119`	`119`	`// denghui replaced at 20221110`
`120`	`120`	`// ModuleBase::GlobalFunc::ZEROS(this->rho[is], this->charge->nrxx);`
`121`		`- setmem_var_op()(this->ctx, this->rho[is], 0, this->charge->nrxx);`
	`121`	`+ setmem_var_op()(this->rho[is], 0, this->charge->nrxx);`
`122`	`122`	`if (get_xc_func_type() == 3)`
`123`	`123`	`{`
`124`	`124`	`// ModuleBase::GlobalFunc::ZEROS(this->charge->kin_r[is], this->charge->nrxx);`
`125`		`- setmem_var_op()(this->ctx, this->kin_r[is], 0, this->charge->nrxx);`
	`125`	`+ setmem_var_op()(this->kin_r[is], 0, this->charge->nrxx);`
`126`	`126`	`}`
`127`	`127`	`if (PARAM.globalv.double_grid \|\| PARAM.globalv.use_uspp)`
`128`	`128`	`{`
`129`		`- setmem_complex_op()(this->ctx, this->rhog[is], 0, this->charge->rhopw->npw);`
	`129`	`+ setmem_complex_op()(this->rhog[is], 0, this->charge->rhopw->npw);`
`130`	`130`	`}`
`131`	`131`	`}`
`132`	`132`
`@@ -244,7 +244,7 @@ void ElecStatePW<T, Device>::rhoBandK(const psi::Psi<T, Device>& psi)`
`244`	`244`	`{`
`245`	`245`	`for (int j = 0; j < 3; j++)`
`246`	`246`	`{`
`247`		`- setmem_complex_op()(this->ctx, this->wfcr, 0, this->charge->nrxx);`
	`247`	`+ setmem_complex_op()(this->wfcr, 0, this->charge->nrxx);`
`248`	`248`
`249`	`249`	`meta_op()(this->ctx,`
`250`	`250`	`ik,`
`@@ -280,7 +280,7 @@ void ElecStatePW<T, Device>::cal_becsum(const psi::Psi<T, Device>& psi)`
`280`	`280`	`resmem_complex_op()(becp, nbands * nkb, "ElecState<PW>::becp");`
`281`	`281`	`const int nh_tot = this->ppcell->nhm * (this->ppcell->nhm + 1) / 2;`
`282`	`282`	`resmem_var_op()(becsum, nh_tot * ucell->nat * PARAM.inp.nspin, "ElecState<PW>::becsum");`
`283`		`- setmem_var_op()(this->ctx, becsum, 0, nh_tot * ucell->nat * PARAM.inp.nspin);`
	`283`	`+ setmem_var_op()(becsum, 0, nh_tot * ucell->nat * PARAM.inp.nspin);`
`284`	`284`
`285`	`285`	`for (int ik = 0; ik < psi.get_nk(); ++ik)`
`286`	`286`	`{`
Original file line number	Diff line number	Diff line change
`@@ -9,7 +9,7 @@ void ElecStatePW<T, Device>::cal_tau(const psi::Psi<T, Device>& psi)`
`9`	`9`	`ModuleBase::TITLE("ElecStatePW", "cal_tau");`
`10`	`10`	`for(int is=0; is<PARAM.inp.nspin; is++)`
`11`	`11`	`{`
`12`		`- setmem_var_op()(this->ctx, this->kin_r[is], 0, this->charge->nrxx);`
	`12`	`+ setmem_var_op()(this->kin_r[is], 0, this->charge->nrxx);`
`13`	`13`	`}`
`14`	`14`
`15`	`15`	`for (int ik = 0; ik < psi.get_nk(); ++ik)`
`@@ -31,7 +31,7 @@ void ElecStatePW<T, Device>::cal_tau(const psi::Psi<T, Device>& psi)`
`31`	`31`	`// kinetic energy density`
`32`	`32`	`for (int j = 0; j < 3; j++)`
`33`	`33`	`{`
`34`		`- setmem_complex_op()(this->ctx, this->wfcr, 0, this->charge->nrxx);`
	`34`	`+ setmem_complex_op()(this->wfcr, 0, this->charge->nrxx);`
`35`	`35`
`36`	`36`	`meta_op()(this->ctx,`
`37`	`37`	`ik,`
Original file line number	Diff line number	Diff line change
`@@ -16,7 +16,7 @@ void ElecStatePW_SDFT<T, Device>::psiToRho(const psi::Psi<T, Device>& psi)`
`16`	`16`	`const int nspin = PARAM.inp.nspin;`
`17`	`17`	`for (int is = 0; is < nspin; is++)`
`18`	`18`	`{`
`19`		`- setmem_var_op()(this->ctx, this->rho[is], 0, this->charge->nrxx);`
	`19`	`+ setmem_var_op()(this->rho[is], 0, this->charge->nrxx);`
`20`	`20`	`}`
`21`	`21`
`22`	`22`	`if (GlobalV::MY_STOGROUP == 0)`