Merge branch 'develop' into fft1

Author: A-006

source/module_base/blas_connector.cpp

@@ -69,17 +69,17 @@ float BlasConnector::dot( const int n, const float *X, const int incX, const flo
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		return sdot_(&n, X, &incX, Y, &incY);
+	}
 	return sdot_(&n, X, &incX, Y, &incY);
 }
-}
 
 double BlasConnector::dot( const int n, const double *X, const int incX, const double *Y, const int incY, base_device::AbacusDevice_t device_type)
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		return ddot_(&n, X, &incX, Y, &incY);
+	}
 	return ddot_(&n, X, &incX, Y, &incY);
 }
-}
 
 // C = a * A.? * B.? + b * C
 void BlasConnector::gemm(const char transa, const char transb, const int m, const int n, const int k,
@@ -196,39 +196,39 @@ float BlasConnector::nrm2( const int n, const float *X, const int incX, base_dev
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		return snrm2_( &n, X, &incX );
+	}
 	return snrm2_( &n, X, &incX );
 }
-}
 
 
 double BlasConnector::nrm2( const int n, const double *X, const int incX, base_device::AbacusDevice_t device_type )
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		return dnrm2_( &n, X, &incX );
+	}
 	return dnrm2_( &n, X, &incX );
 }
-}
 
 
 double BlasConnector::nrm2( const int n, const std::complex<double> *X, const int incX, base_device::AbacusDevice_t device_type )
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		return dznrm2_( &n, X, &incX );
+	}
 	return dznrm2_( &n, X, &incX );
 }
-}
 
 // copies a into b
 void BlasConnector::copy(const long n, const double *a, const int incx, double *b, const int incy, base_device::AbacusDevice_t device_type)
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		dcopy_(&n, a, &incx, b, &incy);
-}
+	}
 }
 
 void BlasConnector::copy(const long n, const std::complex<double> *a, const int incx, std::complex<double> *b, const int incy, base_device::AbacusDevice_t device_type)
 {
 	if (device_type == base_device::AbacusDevice_t::CpuDevice) {
 		zcopy_(&n, a, &incx, b, &incy);
-}
+	}
 }

source/module_elecstate/module_charge/charge_mixing.cpp

@@ -1708,11 +1708,6 @@ bool Charge_Mixing::if_scf_oscillate(const int iteration, const double drho, con
     ModuleBase::TITLE("Charge_Mixing", "if_scf_oscillate");
     ModuleBase::timer::tick("Charge_Mixing", "if_scf_oscillate");
 
-    if(threshold >= 0) // close the function
-    {
-        return false;
-    }
-
     if(this->_drho_history.size() == 0)
     {
         this->_drho_history.resize(PARAM.inp.scf_nmax);
@@ -1721,8 +1716,13 @@ bool Charge_Mixing::if_scf_oscillate(const int iteration, const double drho, con
     // add drho into history
     this->_drho_history[iteration - 1] = drho;
 
+    if(threshold >= 0) // close the function
+    {
+        return false;
+    }
+
     // check if the history is long enough
-    if(iteration < iternum_used)
+    if(iteration < iternum_used + this->mixing_restart_last)
     {
         return false;
     }

source/module_elecstate/module_charge/charge_mixing.h

@@ -102,8 +102,9 @@ class Charge_Mixing
     Base_Mixing::Mixing* get_mixing() const {return mixing;}
 
     // for mixing restart
-    int mixing_restart_step = 0; //which step to restart mixing during SCF
+    int mixing_restart_step = 0; //which step to restart mixing during SCF, always equal to scf_namx except for the mixing restart
     int mixing_restart_count = 0; // the number of restart mixing during SCF. Do not set mixing_restart_count as bool since I want to keep some flexibility in the future
+    int mixing_restart_last = 0; // the label of mixing restart step, store the step number of the last mixing restart
 
     // to calculate the slope of drho curve during SCF, which is used to determine if SCF oscillate
     bool if_scf_oscillate(const int iteration, const double drho, const int iternum_used, const double threshold);

source/module_esolver/esolver_ks.cpp

@@ -681,6 +681,7 @@ void ESolver_KS<T, Device>::iter_finish(const int istep, int& iter)
     // notice for restart
     if (PARAM.inp.mixing_restart > 0 && iter == this->p_chgmix->mixing_restart_step - 1 && iter != PARAM.inp.scf_nmax)
     {
+        this->p_chgmix->mixing_restart_last = iter;
         std::cout << " SCF restart after this step!" << std::endl;
     }
 }

source/module_hsolver/kernels/cuda/math_kernel_op.cu

@@ -12,7 +12,7 @@
 namespace hsolver
 {
 const int warp_size = 32;
-const unsigned int full_mask = 0xffffffff;
+// const unsigned int full_mask = 0xffffffff;
 const int thread_per_block = 256;
 }
 
@@ -65,11 +65,11 @@ void destoryBLAShandle(){
     }
 }
 
-template <typename FPTYPE>
-__forceinline__ __device__ void warp_reduce(FPTYPE& val) {
-    for (int offset = 16; offset > 0; offset >>= 1)
-        val += __shfl_down_sync(full_mask, val, offset);
-}
+// template <typename FPTYPE>
+// __forceinline__ __device__ void warp_reduce(FPTYPE& val) {
+//     for (int offset = 16; offset > 0; offset >>= 1)
+//         val += __shfl_down_sync(full_mask, val, offset);
+// }
 
 template <typename Real>
 __global__ void line_minimize_with_block(