Merge branch 'develop' into kirk0830-kline-fix-1

examples/bsse/water/result.ref

@@ -1 +1,5 @@
--13.3178
+-13.31774791420906
+E_H2O: -466.1219344619547
+E_O: -427.6267373722137
+E_H1: -12.58872471938786
+E_H2: -12.58872445614408

examples/bsse/water/run.sh

@@ -23,14 +23,18 @@ mv OUT.ABACUS/running_scf.log OUT.ABACUS/running_scf_H2.log
 
 rm STRU
 
-E_H2O=`grep GE H2O_scf.output | tail -n 1 | awk '{printf $2}'`
-E_O=`grep GE O_scf.output | tail -n 1 | awk '{printf $2}'`
-E_H1=`grep GE H1_scf.output | tail -n 1 | awk '{printf $2}'`
-E_H2=`grep GE H2_scf.output | tail -n 1 | awk '{printf $2}'`
-result=$(awk 'BEGIN{print "'$E_H2O'" - "'$E_O'" - "'$E_H1'" - "'$E_H2'"}')
+E_H2O=$(grep FINAL_ETOT_IS OUT.ABACUS/running_scf_H2O.log | awk '{printf $2}')
+E_O=$(grep FINAL_ETOT_IS OUT.ABACUS/running_scf_O.log | awk '{printf $2}')
+E_H1=$(grep FINAL_ETOT_IS OUT.ABACUS/running_scf_H1.log | awk '{printf $2}')
+E_H2=$(grep FINAL_ETOT_IS OUT.ABACUS/running_scf_H2.log | awk '{printf $2}')
+result=$(echo "scale=12; ${E_H2O} - ${E_O} - ${E_H1} - ${E_H2}" | bc -l)
 echo $result > result.out
-result_ref=$(cat result.ref)
-difference=$(awk 'BEGIN{print "'$result'" - "'$result_ref'"}')
+echo "E_H2O: $E_H2O" >> result.out
+echo "E_O: $E_O" >> result.out
+echo "E_H1: $E_H1" >> result.out
+echo "E_H2: $E_H2" >> result.out
+result_ref=$(cat result.ref | head -1)
+difference=$(echo "scale=12; $result - $result_ref" | bc -l)
 difference=$(awk -v a=$difference 'BEGIN{if(a<0) print -1*a; else print a}')
 
 if [[ ! -f H2O_scf.output ]] || 

source/module_hamilt_lcao/module_gint/gint_force.h

@@ -192,6 +192,7 @@ void allocateDm(double* matrix_host,
 void para_init(grid_para& para,
                        const int iter_num,
                        const int nbz,
+                       const int pipeline_index,
                        const Grid_Technique& gridt);
 /**
  * @brief frc_strs_iat on host and device Init
@@ -224,6 +225,7 @@ void cal_init(frc_strs_iat& f_s_iat,
 void para_mem_copy(grid_para& para,
                          const Grid_Technique& gridt,
                          const int nbz,
+                         const int pipeline_index,
                          const int atom_num_grid);
 /**
  * @brief Force Stress Force Iat memCpy,from Host to Device

source/module_hamilt_lcao/module_gint/gint_force_gpu.cu

@@ -77,6 +77,7 @@ void gint_fvl_gamma_gpu(hamilt::HContainer<double>* dm,
     const int cuda_block
         = std::min(64, (gridt.psir_size + cuda_threads - 1) / cuda_threads);
     int iter_num = 0;
+    int pipeline_index = 0;
     DensityMat denstiy_mat;
     frc_strs_iat_gbl f_s_iat_dev;
     grid_para para;
@@ -112,9 +113,10 @@ void gint_fvl_gamma_gpu(hamilt::HContainer<double>* dm,
             dim3 grid_dot(cuda_block);
             dim3 block_dot(cuda_threads);
 
-            para_init(para, iter_num, nbz, gridt);
+            pipeline_index = iter_num % gridt.nstreams;
+            para_init(para, iter_num, nbz, pipeline_index,gridt);
             cal_init(f_s_iat,
-                               para.stream_num,
+                               pipeline_index,
                                cuda_block,
                                atom_num_grid,
                                max_size,
@@ -141,19 +143,20 @@ void gint_fvl_gamma_gpu(hamilt::HContainer<double>* dm,
             para_mem_copy(para, 
                                  gridt, 
                                  nbz, 
+                                 pipeline_index,
                                  atom_num_grid);
             cal_mem_cpy(f_s_iat,
                                  gridt,
                                  atom_num_grid,
                                  cuda_block,
-                                 para.stream_num);
-            checkCuda(cudaStreamSynchronize(gridt.streams[para.stream_num]));
+                                 pipeline_index);
+            checkCuda(cudaStreamSynchronize(gridt.streams[pipeline_index]));
             /* cuda stream compute and Multiplication of multinomial matrices */
 
             get_psi_force<<<grid_psi,
                             block_psi,
                             0,
-                            gridt.streams[para.stream_num]>>>(
+                            gridt.streams[pipeline_index]>>>(
                 gridt.ylmcoef_g,
                 dr,
                 gridt.bxyz,
@@ -192,14 +195,14 @@ void gint_fvl_gamma_gpu(hamilt::HContainer<double>* dm,
                                      para.matrix_C_device,
                                      para.ldc_device,
                                      atom_pair_num,
-                                     gridt.streams[para.stream_num],
+                                     gridt.streams[pipeline_index],
                                      nullptr);
             /* force compute in GPU */
             if (isforce){
             dot_product_force<<<grid_dot_force,
                                 block_dot_force,
                                 0,
-                                gridt.streams[para.stream_num]>>>(
+                                gridt.streams[pipeline_index]>>>(
                 para.psir_lx_device,
                 para.psir_ly_device,
                 para.psir_lz_device,
@@ -215,7 +218,7 @@ void gint_fvl_gamma_gpu(hamilt::HContainer<double>* dm,
             dot_product_stress<<<grid_dot,
                                  block_dot,
                                  0,
-                                 gridt.streams[para.stream_num]>>>(
+                                 gridt.streams[pipeline_index]>>>(
                 para.psir_lxx_device,
                 para.psir_lxy_device,
                 para.psir_lxz_device,