Skip to content

Commit ba051ff

Browse files
author
dyzheng
committed
Fix: cuda error of force_op.cu and stress_op.cu
1 parent 9693ecf commit ba051ff

2 files changed

Lines changed: 20 additions & 4 deletions

File tree

source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/force_op.cu

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -429,7 +429,11 @@ __global__ void cal_force_onsite(int wg_nc,
429429
{
430430
iat += atom_na[ii];
431431
sum += atom_na[ii] * atom_nh[ii];
432-
vu += 4 * tlp1_2 * atom_na[ii]; // step for vu
432+
if(orbital_corr[ii] != -1)
433+
{
434+
int size_vu_ii = 4 * (orbital_corr[ii] * 2 + 1) * (orbital_corr[ii] * 2 + 1);
435+
vu += size_vu_ii * atom_na[ii]; // step for vu
436+
}
433437
}
434438

435439
const FPTYPE fac = d_wg[ik * wg_nc + ib] * 2.0 * tpiba;
@@ -496,7 +500,11 @@ __global__ void cal_force_onsite_np1(int wg_nc,
496500
{
497501
iat += atom_na[ii];
498502
sum += atom_na[ii] * atom_nh[ii];
499-
vu += tlp1_2 * atom_na[ii]; // step for vu
503+
if(orbital_corr[ii] != -1)
504+
{
505+
int size_vu_ii = (orbital_corr[ii] * 2 + 1) * (orbital_corr[ii] * 2 + 1);
506+
vu += size_vu_ii * atom_na[ii]; // step for vu
507+
}
500508
}
501509

502510
const FPTYPE fac = d_wg[ik * wg_nc + ib] * 2.0 * tpiba;

source/module_hamilt_pw/hamilt_pwdft/kernels/cuda/stress_op.cu

Lines changed: 10 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -742,7 +742,11 @@ __global__ void cal_stress_onsite(
742742
for (int ii = 0; ii < it; ii++) {
743743
iat += atom_na[ii];
744744
sum += atom_na[ii] * atom_nh[ii];
745-
vu += 4 * tlp1_2 * atom_na[ii];// step for vu
745+
if(orbital_corr[ii] != -1)
746+
{
747+
int size_vu_ii = 4 * (orbital_corr[ii] * 2 + 1) * (orbital_corr[ii] * 2 + 1);
748+
vu += size_vu_ii * atom_na[ii]; // step for vu
749+
}
746750
}
747751

748752
FPTYPE stress_var = 0;
@@ -804,7 +808,11 @@ __global__ void cal_stress_onsite_np1(
804808
for (int ii = 0; ii < it; ii++) {
805809
iat += atom_na[ii];
806810
sum += atom_na[ii] * atom_nh[ii];
807-
vu += tlp1_2 * atom_na[ii];// step for vu
811+
if(orbital_corr[ii] != -1)
812+
{
813+
int size_vu_ii = (orbital_corr[ii] * 2 + 1) * (orbital_corr[ii] * 2 + 1);
814+
vu += size_vu_ii * atom_na[ii]; // step for vu
815+
}
808816
}
809817

810818
FPTYPE stress_var = 0;

0 commit comments

Comments
 (0)