Fix(GPU): fix compilation errors for GPU build with CUDA

- Add unique_cuda_ptr and unique_host_ptr RAII wrappers in cuda.h
  used by cusolver.h for automatic GPU/host memory management
- Fix libxc_abacus.h declarations: add missing hybrid_alpha and
  hse_omega parameters to init_func, v_xc_libxc, v_xc_meta,
  xc_spin_libxc, gcxc_libxc, gcxc_spin_libxc, tau_xc, tau_xc_spin
  to match their implementations
- Fix GPU stress_op.h DeltaSpin: add missing isk parameter to
  match implementation and caller
- Fix GPU force_op.h/cu/hip.cu DeltaSpin: remove isk parameter
  that was absent from CPU version and not passed by caller
- Fix test CMakeLists: add missing source files for linking
  (psi_paging.cpp, spin_constrain.cpp, xc_functional_gga_noncol_sf_builtin.cpp)
- Fix spin_constrain_test: update set_input_parameters call to
  include new sc_acceleration_mode and sc_acceleration_rms_thr params

Verified paged mode (device_memory_mode=paged) produces identical
energy to full_gpu mode on a 2x2x2 k-point PW test case.

Author: dyzheng

source/source_base/module_container/base/macros/cuda.h

@@ -10,6 +10,44 @@
 
 #define THREADS_PER_BLOCK 256
 
+template <typename T>
+class unique_cuda_ptr {
+    T* ptr_;
+public:
+    unique_cuda_ptr() : ptr_(nullptr) {}
+    explicit unique_cuda_ptr(T* ptr) : ptr_(ptr) {}
+    ~unique_cuda_ptr() { if (ptr_) cudaFree(ptr_); }
+    unique_cuda_ptr(const unique_cuda_ptr&) = delete;
+    unique_cuda_ptr& operator=(const unique_cuda_ptr&) = delete;
+    unique_cuda_ptr(unique_cuda_ptr&& other) noexcept : ptr_(other.ptr_) { other.ptr_ = nullptr; }
+    unique_cuda_ptr& operator=(unique_cuda_ptr&& other) noexcept {
+        if (this != &other) { if (ptr_) cudaFree(ptr_); ptr_ = other.ptr_; other.ptr_ = nullptr; }
+        return *this;
+    }
+    T* get() const { return ptr_; }
+    operator T*() const { return ptr_; }
+    void reset(T* ptr = nullptr) { if (ptr_) cudaFree(ptr_); ptr_ = ptr; }
+};
+
+template <typename T>
+class unique_host_ptr {
+    T* ptr_;
+public:
+    unique_host_ptr() : ptr_(nullptr) {}
+    explicit unique_host_ptr(T* ptr) : ptr_(ptr) {}
+    ~unique_host_ptr() { if (ptr_) free(ptr_); }
+    unique_host_ptr(const unique_host_ptr&) = delete;
+    unique_host_ptr& operator=(const unique_host_ptr&) = delete;
+    unique_host_ptr(unique_host_ptr&& other) noexcept : ptr_(other.ptr_) { other.ptr_ = nullptr; }
+    unique_host_ptr& operator=(unique_host_ptr&& other) noexcept {
+        if (this != &other) { if (ptr_) free(ptr_); ptr_ = other.ptr_; other.ptr_ = nullptr; }
+        return *this;
+    }
+    T* get() const { return ptr_; }
+    operator T*() const { return ptr_; }
+    void reset(T* ptr = nullptr) { if (ptr_) free(ptr_); ptr_ = ptr; }
+};
+
 template <typename T>
 struct GetTypeThrust
 {

source/source_estate/test/CMakeLists.txt

@@ -52,6 +52,9 @@ AddTest(
     ../occupy.cpp
     ../module_charge/charge_mpi.cpp
     ../../source_psi/psi.cpp
+    ../../source_psi/psi_paging.cpp
+    ../../source_lcao/module_deltaspin/spin_constrain.cpp
+    ../../source_hamilt/module_xc/xc_functional_gga_noncol_sf_builtin.cpp
     ../../source_base/module_device/memory_op.cpp
 )
 

source/source_hamilt/module_xc/libxc_abacus.h

@@ -44,7 +44,9 @@ namespace XC_Functional_Libxc
      *       followed by a list of parameters.
      */
     extern std::vector<xc_func_type> init_func(const std::vector<int> &func_id, 
-                                               const int xc_polarized);
+                                               const int xc_polarized,
+                                               const double hybrid_alpha = 0.0,
+                                               const double hse_omega = 0.0);
 
     extern void finish_func(std::vector<xc_func_type> &funcs);
 
@@ -59,15 +61,22 @@ namespace XC_Functional_Libxc
         const double &omega, // volume of cell
         const double tpiba,
         const Charge* const chr, // charge density
-        const std::map<int, double>* scaling_factor = nullptr); // added by jghan, 2024-10-10
+        const int nspin_in,
+        const bool domag,
+        const bool domag_z,
+        const std::map<int, double>* scaling_factor = nullptr,
+        const double hybrid_alpha = 0.0,
+        const double hse_omega = 0.0);
 
-    // for mGGA functional
     extern std::tuple<double,double,ModuleBase::matrix,ModuleBase::matrix> v_xc_meta(
         const std::vector<int> &func_id,
         const int &nrxx, // number of real-space grid
         const double &omega, // volume of cell
         const double tpiba,
-        const Charge* const chr);
+        const Charge* const chr,
+        const int nspin,
+        const double hybrid_alpha = 0.0,
+        const double hse_omega = 0.0);
 
 
 //-------------------
@@ -182,7 +191,9 @@ namespace XC_Functional_Libxc
     extern void xc_spin_libxc(
         const std::vector<int> &func_id,
         const double &rhoup, const double &rhodw,
-        double &exc, double &vxcup, double &vxcdw);
+        double &exc, double &vxcup, double &vxcdw,
+        const double hybrid_alpha = 0.0,
+        const double hse_omega = 0.0);
 
 
 //-------------------
@@ -193,14 +204,18 @@ namespace XC_Functional_Libxc
     extern void gcxc_libxc(
         const std::vector<int> &func_id,
         const double &rho, const double &grho,
-        double &sxc, double &v1xc, double &v2xc);
+        double &sxc, double &v1xc, double &v2xc,
+        const double hybrid_alpha = 0.0,
+        const double hse_omega = 0.0);
 
     // the entire GGA functional, for nspin=2 case
     extern void gcxc_spin_libxc(
         const std::vector<int> &func_id,
         const double rhoup, const double rhodw,
         const ModuleBase::Vector3<double> gdr1, const ModuleBase::Vector3<double> gdr2,
-        double &sxc, double &v1xcup, double &v1xcdw, double &v2xcup, double &v2xcdw, double &v2xcud);
+        double &sxc, double &v1xcup, double &v1xcdw, double &v2xcup, double &v2xcdw, double &v2xcud,
+        const double hybrid_alpha = 0.0,
+        const double hse_omega = 0.0);
 
 
 //-------------------
@@ -211,15 +226,19 @@ namespace XC_Functional_Libxc
     extern void tau_xc(
         const std::vector<int> &func_id,
         const double &rho, const double &grho, const double &atau, double &sxc,
-        double &v1xc, double &v2xc, double &v3xc);
+        double &v1xc, double &v2xc, double &v3xc,
+        const double& hybrid_alpha = 0.0,
+        const double& hse_omega = 0.0);
 
     extern void tau_xc_spin(
         const std::vector<int> &func_id,
         double rhoup, double rhodw,
         ModuleBase::Vector3<double> gdr1, ModuleBase::Vector3<double> gdr2,
         double tauup, double taudw,
         double &sxc, double &v1xcup, double &v1xcdw, double &v2xcup, double &v2xcdw, double &v2xcud,
-        double &v3xcup, double &v3xcdw);
+        double &v3xcup, double &v3xcdw,
+        const double& hybrid_alpha = 0.0,
+        const double& hse_omega = 0.0);
 
 } // namespace XC_Functional_Libxc
 

source/source_hamilt/module_xc/test/CMakeLists.txt

@@ -71,6 +71,7 @@ AddTest(
     ../xc_gga_corr.cpp ../xc_lda_corr.cpp ../xc_gga_exch.cpp
     ../xc_lda_exch.cpp ../xc_hcth.cpp
     ../xc_pot.cpp
+    ../xc_functional_gga_noncol_sf_builtin.cpp
     ../libxc_pot.cpp
     ../libxc_tools.cpp
     ../../../source_base/module_external/blas_connector_base.cpp ../../../source_base/module_external/blas_connector_vector.cpp ../../../source_base/module_external/blas_connector_matrix.cpp

source/source_lcao/module_deltaspin/test/spin_constrain_test.cpp

@@ -134,7 +134,7 @@ TYPED_TEST(SpinConstrainTest, SetInputParameters)
     double alpha_trial = 0.01;
     double sccut = 3.0;
     double sc_drop_thr = 1e-3;
-    this->sc.set_input_parameters(sc_thr, nsc, nsc_min, alpha_trial, sccut, sc_drop_thr);
+    this->sc.set_input_parameters(sc_thr, nsc, nsc_min, alpha_trial, sccut, sc_drop_thr, "none", 0.0);
     EXPECT_DOUBLE_EQ(this->sc.get_sc_thr(), sc_thr);
     EXPECT_EQ(this->sc.get_nsc(), nsc);
     EXPECT_EQ(this->sc.get_nsc_min(), nsc_min);

source/source_pw/module_pwdft/kernels/cuda/force_op.cu

@@ -545,16 +545,11 @@ void cal_force_nl_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_dev
                                                                   const FPTYPE& tpiba,
                                                                   const FPTYPE* d_wg,
                                                                   const FPTYPE* lambda,
-                                                                  const int* isk,
                                                                   const std::complex<FPTYPE>* becp,
                                                                   const std::complex<FPTYPE>* dbecp,
                                                                   FPTYPE* force)
 {
     int spin_sign = 1;
-    if (isk != nullptr && npol == 1)
-    {
-        spin_sign = (isk[ik] == 0) ? 1 : -1;
-    }
     if (npol == 1)
     {
         cal_force_onsite<FPTYPE, 1>

source/source_pw/module_pwdft/kernels/force_op.h

@@ -277,7 +277,6 @@ struct cal_force_nl_op<FPTYPE, base_device::DEVICE_GPU>
                     const FPTYPE& tpiba,
                     const FPTYPE* d_wg,
                     const FPTYPE* lambda,
-                    const int* isk,
                     const std::complex<FPTYPE>* becp,
                     const std::complex<FPTYPE>* dbecp,
                     FPTYPE* force);

source/source_pw/module_pwdft/kernels/rocm/force_op.hip.cu

@@ -539,16 +539,11 @@ void cal_force_nl_op<FPTYPE, base_device::DEVICE_GPU>::operator()(const base_dev
                                                                    const FPTYPE& tpiba,
                                                                    const FPTYPE* d_wg,
                                                                    const FPTYPE* lambda,
-                                                                   const int* isk,
                                                                    const std::complex<FPTYPE>* becp,
                                                                    const std::complex<FPTYPE>* dbecp,
                                                                    FPTYPE* force)
 {
     int spin_sign = 1;
-    if (isk != nullptr && npol == 1)
-    {
-        spin_sign = (isk[ik] == 0) ? 1 : -1;
-    }
     if (npol == 1)
     {
         hipLaunchKernelGGL(HIP_KERNEL_NAME(cal_force_onsite<FPTYPE, 1>),

source/source_pw/module_pwdft/kernels/stress_op.h

@@ -358,6 +358,7 @@ struct cal_stress_nl_op<FPTYPE, base_device::DEVICE_GPU>
                     const int* atom_na,
                     const FPTYPE* d_wg,
                     const double* lambda,
+                    const int* isk,
                     const std::complex<FPTYPE>* becp,
                     const std::complex<FPTYPE>* dbecp,
                     FPTYPE* stress);