forked from deepmodeling/abacus-develop
-
Notifications
You must be signed in to change notification settings - Fork 0
Expand file tree
/
Copy pathstress_func_exx.cpp
More file actions
342 lines (306 loc) · 14.1 KB
/
Copy pathstress_func_exx.cpp
File metadata and controls
342 lines (306 loc) · 14.1 KB
1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
142
143
144
145
146
147
148
149
150
151
152
153
154
155
156
157
158
159
160
161
162
163
164
165
166
167
168
169
170
171
172
173
174
175
176
177
178
179
180
181
182
183
184
185
186
187
188
189
190
191
192
193
194
195
196
197
198
199
200
201
202
203
204
205
206
207
208
209
210
211
212
213
214
215
216
217
218
219
220
221
222
223
224
225
226
227
228
229
230
231
232
233
234
235
236
237
238
239
240
241
242
243
244
245
246
247
248
249
250
251
252
253
254
255
256
257
258
259
260
261
262
263
264
265
266
267
268
269
270
271
272
273
274
275
276
277
278
279
280
281
282
283
284
285
286
287
288
289
290
291
292
293
294
295
296
297
298
299
300
301
302
303
304
305
306
307
308
309
310
311
312
313
314
315
316
317
318
319
320
321
322
323
324
325
326
327
328
329
330
331
332
333
334
335
336
337
338
339
340
341
342
#include "stress_pw.h"
#include "global.h"
template <typename FPTYPE, typename Device>
void Stress_PW<FPTYPE, Device>::stress_exx(ModuleBase::matrix& sigma,
const ModuleBase::matrix& wg,
ModulePW::PW_Basis* rhopw,
ModulePW::PW_Basis_K* wfcpw,
const K_Vectors *p_kv,
const psi::Psi <std::complex<FPTYPE>, Device>* d_psi_in, const UnitCell& ucell)
{
double nqs_half1 = 0.5 * p_kv->nmp[0];
double nqs_half2 = 0.5 * p_kv->nmp[1];
double nqs_half3 = 0.5 * p_kv->nmp[2];
bool gamma_extrapolation = PARAM.inp.exx_gamma_extrapolation;
if (!p_kv->get_is_mp())
{
gamma_extrapolation = false;
}
auto isint = [](double x)
{
double epsilon = 1e-6; // this follows the isint judgement in q-e
return std::abs(x - std::round(x)) < epsilon;
};
// T is complex of FPTYPE, if FPTYPE is double, T is std::complex<double>
// but if FPTYPE is std::complex<double>, T is still std::complex<double>
using T = std::complex<FPTYPE>;
using Real = FPTYPE;
using setmem_complex_op = base_device::memory::set_memory_op<T, Device>;
using resmem_complex_op = base_device::memory::resize_memory_op<T, Device>;
using resmem_real_op = base_device::memory::resize_memory_op<Real, Device>;
using delmem_complex_op = base_device::memory::delete_memory_op<T, Device>;
using delmem_real_op = base_device::memory::delete_memory_op<Real, Device>;
using syncmem_complex_op = base_device::memory::synchronize_memory_op<T, Device, Device>;
int nks = wfcpw->nks;
int nqs = wfcpw->nks; // currently q-points downsampling is not supported
double omega = ucell.omega;
double tpiba = ucell.tpiba;
double tpiba2 = ucell.tpiba2;
double omega_inv = 1.0 / omega;
// allocate space
T* psi_nk_real = nullptr;
T* psi_mq_real = nullptr;
T* density_real = nullptr;
T* density_recip = nullptr;
Real* pot = nullptr; // This factor is 2x of the potential in 10.1103/PhysRevB.73.125120
Real* pot_stress = nullptr;
resmem_complex_op()(psi_nk_real, wfcpw->nrxx);
resmem_complex_op()(psi_mq_real, wfcpw->nrxx);
resmem_complex_op()(density_real, rhopw->nrxx);
resmem_complex_op()(density_recip, rhopw->npw);
resmem_real_op()(pot, rhopw->npw * nks * nks);
resmem_real_op()(pot_stress, rhopw->npw * nks * nks);
// prepare the coefficients
double exx_div = 0;
// pasted from op_exx_pw.cpp
{
if (GlobalC::exx_info.info_lip.lambda == 0.0)
{
return;
}
// here we follow the exx_divergence subroutine in q-e (PW/src/exx_base.f90)
double alpha = 10.0 / wfcpw->gk_ecut;
double div = 0;
// this is the \sum_q F(q) part
// temporarily for all k points, should be replaced to q points later
for (int ik = 0; ik < wfcpw->nks; ik++)
{
const ModuleBase::Vector3<double> k_c = wfcpw->kvec_c[ik];
const ModuleBase::Vector3<double> k_d = wfcpw->kvec_d[ik];
#ifdef _OPENMP
#pragma omp parallel for reduction(+:div)
#endif
for (int ig = 0; ig < rhopw->npw; ig++)
{
const ModuleBase::Vector3<double> q_c = k_c + rhopw->gcar[ig];
const ModuleBase::Vector3<double> q_d = k_d + rhopw->gdirect[ig];
double qq = q_c.norm2();
// For gamma_extrapolation (https://doi.org/10.1103/PhysRevB.79.205114)
// 7/8 of the points in the grid are "activated" and 1/8 are disabled.
// grid_factor is designed for the 7/8 of the grid to function like all of the points
double grid_factor = 1;
double extrapolate_grid = 8.0/7.0;
if (gamma_extrapolation)
{
if (isint(q_d[0] * nqs_half1) &&
isint(q_d[1] * nqs_half2) &&
isint(q_d[2] * nqs_half3))
{
grid_factor = 0;
}
else
{
grid_factor = extrapolate_grid;
}
}
if (qq <= 1e-8) continue;
else if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erfc)
{
double hse_omega = GlobalC::exx_info.info_global.hse_omega;
double omega2 = hse_omega * hse_omega;
div += std::exp(-alpha * qq) / qq * (1.0 - std::exp(-qq*tpiba2 / 4.0 / omega2)) * grid_factor;
}
else
{
div += std::exp(-alpha * qq) / qq * grid_factor;
}
}
}
Parallel_Reduce::reduce_pool(div);
// std::cout << "EXX div: " << div << std::endl;
// if (PARAM.inp.dft_functional == "hse")
if (!gamma_extrapolation)
{
if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erfc)
{
double omega = GlobalC::exx_info.info_global.hse_omega;
div += tpiba2 / 4.0 / omega / omega; // compensate for the finite value when qq = 0
}
else
{
div -= alpha;
}
}
div *= ModuleBase::e2 * ModuleBase::FOUR_PI / tpiba2 / wfcpw->nks;
// numerically value the mean value of F(q) in the reciprocal space
// This means we need to calculate the average of F(q) in the first brillouin zone
alpha /= tpiba2;
int nqq = 100000;
double dq = 5.0 / std::sqrt(alpha) / nqq;
double aa = 0.0;
// if (PARAM.inp.dft_functional == "hse")
if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erfc)
{
double hse_omega = GlobalC::exx_info.info_global.hse_omega;
double omega2 = hse_omega * hse_omega;
#ifdef _OPENMP
#pragma omp parallel for reduction(+:aa)
#endif
for (int i = 0; i < nqq; i++)
{
double q = dq * (i+0.5);
aa -= exp(-alpha * q * q) * exp(-q*q / 4.0 / omega2) * dq;
}
}
aa *= 8 / ModuleBase::FOUR_PI;
aa += 1.0 / std::sqrt(alpha * ModuleBase::PI);
div -= ModuleBase::e2 * omega * aa;
exx_div = div * wfcpw->nks;
// std::cout << "EXX divergence: " << exx_div << std::endl;
}
// prepare for the potential
for (int ik = 0; ik < nks; ik++)
{
for (int iq = 0; iq < nks; iq++)
{
const ModuleBase::Vector3<double> k_c = wfcpw->kvec_c[ik];
const ModuleBase::Vector3<double> k_d = wfcpw->kvec_d[ik];
const ModuleBase::Vector3<double> q_c = wfcpw->kvec_c[iq];
const ModuleBase::Vector3<double> q_d = wfcpw->kvec_d[iq];
#ifdef _OPENMP
#pragma omp parallel for schedule(static)
#endif
for (int ig = 0; ig < rhopw->npw; ig++)
{
const ModuleBase::Vector3<double> g_d = rhopw->gdirect[ig];
const ModuleBase::Vector3<double> kqg_d = k_d - q_d + g_d;
// For gamma_extrapolation (https://doi.org/10.1103/PhysRevB.79.205114)
// 7/8 of the points in the grid are "activated" and 1/8 are disabled.
// grid_factor is designed for the 7/8 of the grid to function like all of the points
Real grid_factor = 1;
if (gamma_extrapolation)
{
double extrapolate_grid = 8.0/7.0;
if (isint(kqg_d[0] * nqs_half1) &&
isint(kqg_d[1] * nqs_half2) &&
isint(kqg_d[2] * nqs_half3))
{
grid_factor = 0;
}
else
{
grid_factor = extrapolate_grid;
}
}
const int ig_kq = ik * nks * rhopw->npw + iq * rhopw->npw + ig;
Real gg = (k_c - q_c + rhopw->gcar[ig]).norm2() * tpiba2;
Real hse_omega2 = GlobalC::exx_info.info_global.hse_omega * GlobalC::exx_info.info_global.hse_omega;
// if (kqgcar2 > 1e-12) // vasp uses 1/40 of the smallest (k spacing)**2
if (gg >= 1e-8)
{
Real fac = -ModuleBase::FOUR_PI * ModuleBase::e2 / gg;
// if (PARAM.inp.dft_functional == "hse")
if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erfc)
{
pot[ig_kq] = fac * (1.0 - std::exp(-gg / 4.0 / hse_omega2)) * grid_factor;
pot_stress[ig_kq] = (1.0 - (1.0 + gg / 4.0 / hse_omega2) * std::exp(-gg / 4.0 / hse_omega2)) / (1.0 - std::exp(-gg / 4.0 / hse_omega2)) / gg;
}
else if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erf)
{
ModuleBase::WARNING("Stress_PW", "Stress for Erf is not implemented yet");
pot[ig_kq] = fac * grid_factor;
pot_stress[ig_kq] = 1.0 / gg;
}
else if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Hf)
{
pot[ig_kq] = fac * grid_factor;
pot_stress[ig_kq] = 1.0 / gg;
}
}
// }
else
{
// if (PARAM.inp.dft_functional == "hse")
if (GlobalC::exx_info.info_global.ccp_type == Conv_Coulomb_Pot_K::Ccp_Type::Erfc && !gamma_extrapolation)
{
pot[ig_kq] = - ModuleBase::PI * ModuleBase::e2 / hse_omega2; // maybe we should add a exx_div here, but q-e does not do that
pot_stress[ig_kq] = 1 / 4.0 / hse_omega2;
}
else
{
pot[ig_kq] = exx_div;
pot_stress[ig_kq] = 0;
}
}
// assert(is_finite(density_recip[ig]));
}
}
}
// calculate the stress
// for nk, mq
for (int ik = 0; ik < nks; ik++)
{
for (int nband = 0; nband < d_psi_in->get_nbands(); nband++)
{
if (wg(ik, nband) < 1e-12) continue;
// psi_nk in real space
d_psi_in->fix_kb(ik, nband);
T* psi_nk = d_psi_in->get_pointer();
wfcpw->recip_to_real<std::complex<FPTYPE>,Device>(psi_nk, psi_nk_real, ik);
for (int iq = 0; iq < nqs; iq++)
{
for (int mband = 0; mband < d_psi_in->get_nbands(); mband++)
{
// psi_mq in real space
d_psi_in->fix_kb(iq, mband);
T* psi_mq = d_psi_in->get_pointer();
wfcpw->recip_to_real<std::complex<FPTYPE>,Device>(psi_mq, psi_mq_real, iq);
// overlap density in real space
setmem_complex_op()(density_real, 0.0, rhopw->nrxx);
for (int ig = 0; ig < rhopw->nrxx; ig++)
{
density_real[ig] = psi_nk_real[ig] * std::conj(psi_mq_real[ig]) * omega_inv;
}
// density in reciprocal space
rhopw->real2recip(density_real, density_recip);
// really calculate the stress
// for alpha beta
for (int alpha = 0; alpha < 3; alpha++)
{
for (int beta = alpha; beta < 3; beta++)
{
int delta_ab = (alpha == beta) ? 1 : 0;
double sigma_ab_loc = 0.0;
#ifdef _OPENMP
#pragma omp parallel for schedule(static) reduction(+:sigma_ab_loc)
#endif
for (int ig = 0; ig < rhopw->npw; ig++)
{
const ModuleBase::Vector3<double> kqg = wfcpw->kvec_c[ik] - wfcpw->kvec_c[iq] + rhopw->gcar[ig];
double kqg_alpha = kqg[alpha] * tpiba;
double kqg_beta = kqg[beta] * tpiba;
// equation 10 of 10.1103/PhysRevB.73.125120
double density_recip2 = std::real(density_recip[ig] * std::conj(density_recip[ig]));
const int idx = ig + iq * rhopw->npw + ik * rhopw->npw * nqs;
double pot_local = pot[idx];
double pot_stress_local = pot_stress[idx];
sigma_ab_loc += density_recip2 * pot_local * (kqg_alpha * kqg_beta * pot_stress_local - delta_ab) ;
}
// 0.5 in the following line is caused by 2x in the pot
sigma(alpha, beta) -= GlobalC::exx_info.info_global.hybrid_alpha
* 0.25 * sigma_ab_loc
* wg(ik, nband) * wg(iq, mband) / nqs / p_kv->wk[ik];
}
}
}
}
}
}
for (int l = 0; l < 3; l++)
{
for (int m = l + 1; m < 3; m++)
{
sigma(m, l) = sigma(l, m);
}
}
Parallel_Reduce::reduce_all(sigma.c, sigma.nr * sigma.nc);
delmem_complex_op()(psi_nk_real);
delmem_complex_op()(psi_mq_real);
delmem_complex_op()(density_real);
delmem_complex_op()(density_recip);
delmem_real_op()(pot);
}
template class Stress_PW<double, base_device::DEVICE_CPU>;
#if ((defined __CUDA) || (defined __ROCM))
template class Stress_PW<double, base_device::DEVICE_GPU>;
#endif