Remove ctx in axpy_op

Critsium-xy · Critsium-xy · commit 9db93de477bf · 2025-02-12T13:11:44.000+08:00
diff --git a/source/module_base/kernels/cuda/math_kernel_op.cu b/source/module_base/kernels/cuda/math_kernel_op.cu
@@ -718,8 +718,7 @@ cublasOperation_t judge_trans_op(bool is_complex, const char& trans, const char*
 }
 
 template <>
-void gemv_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                          const char& trans,
+void gemv_op<double, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                           const int& m,
                                                           const int& n,
                                                           const double* alpha,
@@ -736,8 +735,7 @@ void gemv_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEV
 }
 
 template <>
-void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                                       const char& trans,
+void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                                        const int& m,
                                                                        const int& n,
                                                                        const std::complex<float>* alpha_in,
@@ -756,8 +754,7 @@ void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const bas
 }
 
 template <>
-void gemv_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                                        const char& trans,
+void gemv_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                                         const int& m,
                                                                         const int& n,
                                                                         const std::complex<double>* alpha_in,
diff --git a/source/module_base/kernels/math_kernel_op.cpp b/source/module_base/kernels/math_kernel_op.cpp
@@ -231,8 +231,7 @@ struct scal_op<FPTYPE, base_device::DEVICE_CPU>
 template <typename T>
 struct gemv_op<T, base_device::DEVICE_CPU>
 {
-    void operator()(const base_device::DEVICE_CPU* d,
-                    const char& trans,
+    void operator()(const char& trans,
                     const int& m,
                     const int& n,
                     const T* alpha,
diff --git a/source/module_base/kernels/math_kernel_op.h b/source/module_base/kernels/math_kernel_op.h
@@ -208,7 +208,6 @@ template <typename T, typename Device> struct gemv_op {
   /// @brief y = alpha * op(A) * x + beta * y
   ///
   /// Input Parameters
-  /// \param d : the type of computing device
   /// \param trans : whether to transpose A
   /// \param m : first dimension of matrix
   /// \param n : second dimension of matrix
@@ -223,7 +222,7 @@ template <typename T, typename Device> struct gemv_op {
   ///
   /// Output Parameters
   /// \param Y : output array Y
-  void operator()(const Device *d, const char &trans, const int &m,
+  void operator()(const char &trans, const int &m,
                   const int &n, const T *alpha, const T *A, const int &lda,
                   const T *X, const int &incx, const T *beta, T *Y,
                   const int &incy);
diff --git a/source/module_base/kernels/rocm/math_kernel_op.hip.cu b/source/module_base/kernels/rocm/math_kernel_op.hip.cu
@@ -642,8 +642,7 @@ hipblasOperation_t judge_trans_op(bool is_complex, const char& trans, const char
 }
 
 template <>
-void gemv_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                          const char& trans,
+void gemv_op<double, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                           const int& m,
                                                           const int& n,
                                                           const double* alpha,
@@ -660,8 +659,7 @@ void gemv_op<double, base_device::DEVICE_GPU>::operator()(const base_device::DEV
 }
 
 template <>
-void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                                       const char& trans,
+void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                                        const int& m,
                                                                        const int& n,
                                                                        const std::complex<float>* alpha,
@@ -678,8 +676,7 @@ void gemv_op<std::complex<float>, base_device::DEVICE_GPU>::operator()(const bas
 }
 
 template <>
-void gemv_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const base_device::DEVICE_GPU* d,
-                                                                        const char& trans,
+void gemv_op<std::complex<double>, base_device::DEVICE_GPU>::operator()(const char& trans,
                                                                         const int& m,
                                                                         const int& n,
                                                                         const std::complex<double>* alpha,
diff --git a/source/module_base/kernels/test/math_kernel_test.cpp b/source/module_base/kernels/test/math_kernel_test.cpp
@@ -332,8 +332,7 @@ TEST_F(TestModuleHsolverMathKernel, scal_op_cpu)
 
 TEST_F(TestModuleHsolverMathKernel, gemv_op_cpu)
 {
-    gemv_op_cpu()(cpu_ctx,
-                  'C',
+    gemv_op_cpu()('C',
                   2,
                   3,
                   &ModuleBase::ONE,
@@ -598,7 +597,7 @@ TEST_F(TestModuleHsolverMathKernel, gemv_op_gpu)
 
     // run
     ModuleBase::createGpuBlasHandle();
-    gemv_op_gpu()(gpu_ctx, 'C', 2, 3, &ModuleBase::ONE, A_gemv_dev, 2, X_gemv_dev, 1, &ModuleBase::ONE, Y_gemv_dev, 1);
+    gemv_op_gpu()('C', 2, 3, &ModuleBase::ONE, A_gemv_dev, 2, X_gemv_dev, 1, &ModuleBase::ONE, Y_gemv_dev, 1);
     ModuleBase::destoryBLAShandle();
     // syn the output data in GPU to CPU
     synchronize_memory_op_gpu()(Y_gemv.data(), Y_gemv_dev, Y_gemv.size());
diff --git a/source/module_elecstate/elecstate_pw.cpp b/source/module_elecstate/elecstate_pw.cpp
@@ -302,8 +302,7 @@ void ElecStatePW<T, Device>::cal_becsum(const psi::Psi<T, Device>& psi)
         if (nbands == 1)
         {
             int inc = 1;
-            gemv_op()(this->ctx,
-                      transa,
+            gemv_op()(transa,
                       npw,
                       this->ppcell->nkb,
                       &one,
diff --git a/source/module_hamilt_pw/hamilt_pwdft/hamilt_pw.cpp b/source/module_hamilt_pw/hamilt_pwdft/hamilt_pw.cpp
@@ -260,8 +260,7 @@ void HamiltPW<T, Device>::sPsi(const T* psi_in, // psi
             if (nbands == 1)
             {
                 int inc = 1;
-                gemv_op()(this->ctx,
-                          transa,
+                gemv_op()(transa,
                           npw,
                           this->ppcell->nkb,
                           &one,
@@ -351,8 +350,7 @@ void HamiltPW<T, Device>::sPsi(const T* psi_in, // psi
             if (nbands == 1)
             {
                 const int inc = 1;
-                gemv_op()(this->ctx,
-                          transa,
+                gemv_op()(transa,
                           npw,
                           this->ppcell->nkb,
                           &one,
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
@@ -170,7 +170,6 @@ void Nonlocal<OperatorPW<T, Device>>::add_nonlocal_pp(T *hpsi_in, const T *becp,
         // denghui replace 2022-10-20
         // <<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<
         gemv_op()(
-            this->ctx,
             transa,
             this->npw,
             this->ppcell->nkb,
@@ -246,7 +245,6 @@ void Nonlocal<OperatorPW<T, Device>>::act(
                 // denghui replace 2022-10-20
                 // <<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<
                 gemv_op()(
-                    this->ctx,
                     transa,
                     this->npw,
                     nkb,
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_che.h b/source/module_hamilt_pw/hamilt_stodft/sto_che.h
@@ -51,7 +51,7 @@ REAL vTMv(const REAL* v, const REAL* M, const int n)
     const REAL zero = 0;
     REAL* y = nullptr;
     base_device::memory::resize_memory_op<REAL, Device>()(y, n);
-    ModuleBase::gemv_op<REAL, Device>()(ctx, normal, n, n, &one, M, n, v, inc, &zero, y, inc);
+    ModuleBase::gemv_op<REAL, Device>()(normal, n, n, &one, M, n, v, inc, &zero, y, inc);
     REAL result = 0;
     REAL* dot_device = nullptr;
     base_device::memory::resize_memory_op<REAL, Device>()(dot_device, 1);
diff --git a/source/module_hamilt_pw/hamilt_stodft/sto_iter.cpp b/source/module_hamilt_pw/hamilt_stodft/sto_iter.cpp
@@ -769,7 +769,7 @@ void Stochastic_Iter<T, Device>::calTnchi_ik(const int& ik, Stochastic_WF<T, Dev
         T* coef_real = nullptr;
         resmem_complex_op()(coef_real, N);
         castmem_d2z_op()(coef_real, p_che->coef_real, p_che->norder);
-        gemv_op()(this->ctx, transa, M, N, &one, stowf.chiallorder[ik].get_pointer(), LDA, coef_real, inc, &zero, out, inc);
+        gemv_op()(transa, M, N, &one, stowf.chiallorder[ik].get_pointer(), LDA, coef_real, inc, &zero, out, inc);
         // zgemv_(&transa, &M, &N, &one, stowf.chiallorder[ik].get_pointer(), &LDA, coef_real, &inc, &zero, out, &inc);
         delmem_complex_op()(coef_real);
     }
diff --git a/source/module_hsolver/diago_cg.cpp b/source/module_hsolver/diago_cg.cpp
@@ -263,8 +263,7 @@ void DiagoCG<T, Device>::orth_grad(const ct::Tensor& psi,
                                    ct::Tensor& lagrange)
 {
     this->spsi_func_(grad, scg); // scg = S|grad>
-    ModuleBase::gemv_op<T, Device>()(ctx_,
-                                     'C',
+    ModuleBase::gemv_op<T, Device>()('C',
                                      this->n_basis_,
                                      m,
                                      this->one_,
@@ -281,8 +280,7 @@ void DiagoCG<T, Device>::orth_grad(const ct::Tensor& psi,
     // (3) orthogonal |g> and |scg> to all states (0~m-1)
     //<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<
     // haozhihan replace 2022-10-07
-    ModuleBase::gemv_op<T, Device>()(ctx_,
-                                     'N',
+    ModuleBase::gemv_op<T, Device>()('N',
                                      this->n_basis_,
                                      m,
                                      this->neg_one_,
@@ -294,8 +292,7 @@ void DiagoCG<T, Device>::orth_grad(const ct::Tensor& psi,
                                      grad.data<T>(),
                                      1);
 
-    ModuleBase::gemv_op<T, Device>()(ctx_,
-                                     'N',
+    ModuleBase::gemv_op<T, Device>()('N',
                                      this->n_basis_,
                                      m,
                                      this->neg_one_,
@@ -491,8 +488,7 @@ void DiagoCG<T, Device>::schmit_orth(const int& m, const ct::Tensor& psi, const
     //<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<
     // haozhihan replace 2022-10-6
     int inc = 1;
-    ModuleBase::gemv_op<T, Device>()(ctx_,
-                                     'C',
+    ModuleBase::gemv_op<T, Device>()('C',
                                      this->n_basis_,
                                      m + 1,
                                      this->one_,
@@ -509,8 +505,7 @@ void DiagoCG<T, Device>::schmit_orth(const int& m, const ct::Tensor& psi, const
 
     //<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<<
     // haozhihan replace 2022-10-6
-    ModuleBase::gemv_op<T, Device>()(ctx_,
-                                     'N',
+    ModuleBase::gemv_op<T, Device>()('N',
                                      this->n_basis_,
                                      m,
                                      this->neg_one_,
diff --git a/source/module_hsolver/diago_david.cpp b/source/module_hsolver/diago_david.cpp
@@ -912,8 +912,7 @@ void DiagoDavid<T, Device>::SchmidtOrth(const int& dim,
     // calculate other lagranges for this band
     // lagrange_m[m - mv_size + 1]
     // = basis[m - mv_size + 1]' * spsi[m]
-    ModuleBase::gemv_op<T, Device>()(this->ctx,
-                                     'C',
+    ModuleBase::gemv_op<T, Device>()('C',
                                      dim,
                                      mv_size,
                                      this->one,
@@ -935,8 +934,7 @@ void DiagoDavid<T, Device>::SchmidtOrth(const int& dim,
 
     // / psi_m = psi_m - \sum_{i < m} \langle psi(i)|S|psi(m) \rangle psi(i)
     // psi_m = psi_m - basis * lagrange_m
-    ModuleBase::gemv_op<T, Device>()(this->ctx,
-                                     'N',
+    ModuleBase::gemv_op<T, Device>()('N',
                                      dim,
                                      m,
                                      this->neg_one,
diff --git a/source/module_hsolver/diago_iter_assist.cpp b/source/module_hsolver/diago_iter_assist.cpp
@@ -222,7 +222,7 @@ void DiagoIterAssist<T, Device>::diagH_subspace_init(hamilt::Hamilt<T, Device>*
             pHamilt->ops->hPsi(hpsi_in);
 
             // calculate the related elements in hcc <Psi|H|Psi>
-            ModuleBase::gemv_op<T, Device>()(ctx, 'C', psi_nc, nstart, &one, psi, psi_nc, hpsi, 1, &zero, hcc + i * nstart, 1);
+            ModuleBase::gemv_op<T, Device>()('C', psi_nc, nstart, &one, psi, psi_nc, hpsi, 1, &zero, hcc + i * nstart, 1);
         }
 
         T* spsi = temp;
@@ -232,8 +232,7 @@ void DiagoIterAssist<T, Device>::diagH_subspace_init(hamilt::Hamilt<T, Device>*
             syncmem_complex_op()(ppsi, psi + i * psi_nc, psi_nc);
             pHamilt->sPsi(ppsi, spsi, dmin, dmin, 1);
 
-            ModuleBase::gemv_op<T, Device>()(ctx,
-                                             'C',
+            ModuleBase::gemv_op<T, Device>()('C',
                                              psi_nc,
                                              nstart,
                                              &one,

Original file line number	Diff line number	Diff line change
`@@ -769,7 +769,7 @@ void Stochastic_Iter<T, Device>::calTnchi_ik(const int& ik, Stochastic_WF<T, Dev`
`769`	`769`	`T* coef_real = nullptr;`
`770`	`770`	`resmem_complex_op()(coef_real, N);`
`771`	`771`	`castmem_d2z_op()(coef_real, p_che->coef_real, p_che->norder);`
`772`		`- gemv_op()(this->ctx, transa, M, N, &one, stowf.chiallorder[ik].get_pointer(), LDA, coef_real, inc, &zero, out, inc);`
	`772`	`+ gemv_op()(transa, M, N, &one, stowf.chiallorder[ik].get_pointer(), LDA, coef_real, inc, &zero, out, inc);`
`773`	`773`	`// zgemv_(&transa, &M, &N, &one, stowf.chiallorder[ik].get_pointer(), &LDA, coef_real, &inc, &zero, out, &inc);`
`774`	`774`	`delmem_complex_op()(coef_real);`
`775`	`775`	`}`