Skip to content

Commit

Permalink
Merge branch 'develop' into lowf-refactor-1
Browse files Browse the repository at this point in the history
  • Loading branch information
kirk0830 authored May 30, 2024
2 parents d40f130 + ec9cc5e commit 4bfff38
Show file tree
Hide file tree
Showing 6 changed files with 394 additions and 380 deletions.
6 changes: 5 additions & 1 deletion examples/bsse/water/result.ref
Original file line number Diff line number Diff line change
@@ -1 +1,5 @@
-13.3178
-13.31774791420906
E_H2O: -466.1219344619547
E_O: -427.6267373722137
E_H1: -12.58872471938786
E_H2: -12.58872445614408
18 changes: 11 additions & 7 deletions examples/bsse/water/run.sh
Original file line number Diff line number Diff line change
Expand Up @@ -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 ]] ||
Expand Down
2 changes: 2 additions & 0 deletions source/module_hamilt_lcao/module_gint/gint_force.h
Original file line number Diff line number Diff line change
Expand Up @@ -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
Expand Down Expand Up @@ -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
Expand Down
19 changes: 11 additions & 8 deletions source/module_hamilt_lcao/module_gint/gint_force_gpu.cu
Original file line number Diff line number Diff line change
Expand Up @@ -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;
Expand Down Expand Up @@ -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,
Expand All @@ -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,
Expand Down Expand Up @@ -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,
Expand All @@ -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,
Expand Down
Loading

0 comments on commit 4bfff38

Please sign in to comment.