Skip to content

Commit 23d9da3

Browse files
author
dyzheng
committed
Fix: force and stress error of DFTU and DeltaSpin for y-axis
1 parent 94099e7 commit 23d9da3

6 files changed

Lines changed: 11 additions & 11 deletions

File tree

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

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -457,7 +457,7 @@ __global__ void cal_force_onsite(int wg_nc,
457457
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb0]) * becp[inkb2 + nkb];
458458
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb0 + nkb]) * becp[inkb2];
459459
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb0 + nkb]) * becp[inkb2 + nkb];
460-
const FPTYPE tmp = -fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
460+
const FPTYPE tmp = -fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
461461
atomicAdd(force + iat * forcenl_nc + ipol, tmp);
462462
}
463463
}

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

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -767,7 +767,7 @@ __global__ void cal_stress_onsite(
767767
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb1]) * becp[inkb2 + nkb];
768768
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb1 + nkb]) * becp[inkb2];
769769
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb1 + nkb]) * becp[inkb2 + nkb];
770-
stress_var -= fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
770+
stress_var -= fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
771771
}
772772
++iat;
773773
sum+=nprojs;
@@ -882,7 +882,7 @@ __global__ void cal_stress_onsite(
882882
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb]) * becp[inkb + nkb];
883883
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb + nkb]) * becp[inkb];
884884
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb + nkb]) * becp[inkb + nkb];
885-
stress_var -= fac * (coefficients0 * dbb0 + coefficients1 * dbb1 + coefficients2 * dbb2 + coefficients3 * dbb3).real();
885+
stress_var -= fac * (coefficients0 * dbb0 + coefficients1 * dbb2 + coefficients2 * dbb1 + coefficients3 * dbb3).real();
886886
}
887887
++iat;
888888
sum+=nprojs;

source/module_hamilt_pw/hamilt_pwdft/kernels/force_op.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -216,7 +216,7 @@ struct cal_force_nl_op<FPTYPE, base_device::DEVICE_CPU>
216216
const std::complex<FPTYPE> dbb2 = conj(dbecp[index0 + nkb]) * becp[index1];
217217
const std::complex<FPTYPE> dbb3 = conj(dbecp[index0 + nkb]) * becp[index1 + nkb];
218218

219-
local_force[ipol] -= fac * (ps0 * dbb0 + ps1 * dbb1 + ps2 * dbb2 + ps3 * dbb3).real();
219+
local_force[ipol] -= fac * (ps0 * dbb0 + ps1 * dbb2 + ps2 * dbb1 + ps3 * dbb3).real();
220220
}
221221
}
222222
}
@@ -317,7 +317,7 @@ struct cal_force_nl_op<FPTYPE, base_device::DEVICE_CPU>
317317
const std::complex<FPTYPE> dbb2 = conj(dbecp[index0 + nkb]) * becp[index1];
318318
const std::complex<FPTYPE> dbb3 = conj(dbecp[index0 + nkb]) * becp[index1 + nkb];
319319

320-
local_force[iforce] -= fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
320+
local_force[iforce] -= fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
321321
}
322322
}
323323
else if(npol == 1)

source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/force_op.hip.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -473,7 +473,7 @@ __global__ void cal_force_onsite(int wg_nc,
473473
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb0]) * becp[inkb2 + nkb];
474474
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb0 + nkb]) * becp[inkb2];
475475
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb0 + nkb]) * becp[inkb2 + nkb];
476-
const FPTYPE tmp = -fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
476+
const FPTYPE tmp = -fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
477477
atomicAdd(force + iat * forcenl_nc + ipol, tmp);
478478
}
479479
}
@@ -596,7 +596,7 @@ __global__ void cal_force_onsite(int wg_nc,
596596
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb0 + nkb]) * becp[inkb + nkb];
597597
const FPTYPE tmp
598598
= -fac
599-
* (coefficients0 * dbb0 + coefficients1 * dbb1 + coefficients2 * dbb2 + coefficients3 * dbb3)
599+
* (coefficients0 * dbb0 + coefficients1 * dbb2 + coefficients2 * dbb1 + coefficients3 * dbb3)
600600
.real();
601601
atomicAdd(force + iat * forcenl_nc + ipol, tmp);
602602
}

source/module_hamilt_pw/hamilt_pwdft/kernels/rocm/stress_op.hip.cu

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -767,7 +767,7 @@ __global__ void cal_stress_onsite(
767767
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb1]) * becp[inkb2 + nkb];
768768
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb1 + nkb]) * becp[inkb2];
769769
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb1 + nkb]) * becp[inkb2 + nkb];
770-
stress_var -= fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
770+
stress_var -= fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
771771
}
772772
++iat;
773773
sum+=nprojs;
@@ -882,7 +882,7 @@ __global__ void cal_stress_onsite(
882882
const thrust::complex<FPTYPE> dbb1 = conj(dbecp[inkb]) * becp[inkb + nkb];
883883
const thrust::complex<FPTYPE> dbb2 = conj(dbecp[inkb + nkb]) * becp[inkb];
884884
const thrust::complex<FPTYPE> dbb3 = conj(dbecp[inkb + nkb]) * becp[inkb + nkb];
885-
stress_var -= fac * (coefficients0 * dbb0 + coefficients1 * dbb1 + coefficients2 * dbb2 + coefficients3 * dbb3).real();
885+
stress_var -= fac * (coefficients0 * dbb0 + coefficients1 * dbb2 + coefficients2 * dbb1 + coefficients3 * dbb3).real();
886886
}
887887
++iat;
888888
sum+=nprojs;

source/module_hamilt_pw/hamilt_pwdft/kernels/stress_op.cpp

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -283,7 +283,7 @@ struct cal_stress_nl_op<FPTYPE, base_device::DEVICE_CPU>
283283
const std::complex<FPTYPE> dbb1 = conj(dbecp[inkb1]) * becp[nkb + inkb2];
284284
const std::complex<FPTYPE> dbb2 = conj(dbecp[nkb + inkb1]) * becp[inkb2];
285285
const std::complex<FPTYPE> dbb3 = conj(dbecp[nkb + inkb1]) * becp[nkb + inkb2];
286-
local_stress -= fac * (ps[0] * dbb0 + ps[1] * dbb1 + ps[2] * dbb2 + ps[3] * dbb3).real();
286+
local_stress -= fac * (ps[0] * dbb0 + ps[1] * dbb2 + ps[2] * dbb1 + ps[3] * dbb3).real();
287287
}
288288
} // end ip
289289
break;
@@ -339,7 +339,7 @@ struct cal_stress_nl_op<FPTYPE, base_device::DEVICE_CPU>
339339
const std::complex<FPTYPE> dbb1 = conj(dbecp[inkb1]) * becp[nkb + inkb1];
340340
const std::complex<FPTYPE> dbb2 = conj(dbecp[nkb + inkb1]) * becp[inkb1];
341341
const std::complex<FPTYPE> dbb3 = conj(dbecp[nkb + inkb1]) * becp[nkb + inkb1];
342-
local_stress -= fac * (coefficients0 * dbb0 + coefficients1 * dbb1 + coefficients2 * dbb2 + coefficients3 * dbb3).real();
342+
local_stress -= fac * (coefficients0 * dbb0 + coefficients1 * dbb2 + coefficients2 * dbb1 + coefficients3 * dbb3).real();
343343
} // end ip
344344
}// ib
345345
}

0 commit comments

Comments
 (0)