Skip to content

Commit 174bbe4

Browse files
authored
Merge pull request #203 from momokchung/exchpol
Exchange Polarization
2 parents 6638f9f + 2503915 commit 174bbe4

38 files changed

+3076
-136
lines changed

.clang-format

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -10,6 +10,7 @@ BreakConstructorInitializers: BeforeComma
1010
PointerAlignment: Left
1111
AllowShortBlocksOnASingleLine: Always
1212
AllowShortFunctionsOnASingleLine: Empty
13+
AllowShortIfStatementsOnASingleLine: WithoutElse
1314

1415
MaxEmptyLinesToKeep: 1
1516
AccessModifierOffset: -3

ext/ext/yaml/alterpol_cu1.yaml

Lines changed: 89 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,89 @@
1+
KERNEL_NAME: alterpol_cu1
2+
3+
CUT_DISTANCE:
4+
- 'cut'
5+
6+
OFF_DISTANCE:
7+
- 'off'
8+
9+
EXCLUDE_INFO:
10+
- 'dinfo'
11+
12+
SCALE_1X_TYPE: real3_const_array,1
13+
14+
EXTRA_PARAMS: |
15+
, real (*restrict polscale)[9]
16+
, const real* restrict kpep, const real* restrict prepep
17+
, const real* restrict dmppep, const int* restrict lpep
18+
, ExpolScr scrtyp
19+
20+
I_POSITION:
21+
- def: shared real xi from:x
22+
- def: shared real yi from:y
23+
- def: shared real zi from:z
24+
K_POSITION:
25+
- def: register real xk from:x
26+
- def: register real yk from:y
27+
- def: register real zk from:z
28+
29+
I_FORCE:
30+
- def: shared real psci00 addto:polscale,0
31+
- def: shared real psci01 addto:polscale,1
32+
- def: shared real psci02 addto:polscale,2
33+
- def: shared real psci10 addto:polscale,3
34+
- def: shared real psci11 addto:polscale,4
35+
- def: shared real psci12 addto:polscale,5
36+
- def: shared real psci20 addto:polscale,6
37+
- def: shared real psci21 addto:polscale,7
38+
- def: shared real psci22 addto:polscale,8
39+
K_FORCE:
40+
- def: register real psck00 addto:polscale,0
41+
- def: register real psck01 addto:polscale,1
42+
- def: register real psck02 addto:polscale,2
43+
- def: register real psck10 addto:polscale,3
44+
- def: register real psck11 addto:polscale,4
45+
- def: register real psck12 addto:polscale,5
46+
- def: register real psck20 addto:polscale,6
47+
- def: register real psck21 addto:polscale,7
48+
- def: register real psck22 addto:polscale,8
49+
50+
I_VARIABLES:
51+
- def: shared real springi from:kpep
52+
- def: shared real sizi from:prepep
53+
- def: shared real alphai from:dmppep
54+
- def: shared int epli from:lpep
55+
K_VARIABLES:
56+
- def: register real springk from:kpep
57+
- def: register real sizk from:prepep
58+
- def: register real alphak from:dmppep
59+
- def: register int eplk from:lpep
60+
61+
FULL_PAIRWISE_INTERACTION: |
62+
real xr = xk - @xi@;
63+
real yr = yk - @yi@;
64+
real zr = zk - @zi@;
65+
real r2 = image2(xr, yr, zr);
66+
if ((eplk or @epli@) and r2 <= off * off and incl) {
67+
real r = REAL_SQRT(r2);
68+
real ks2i[3][3], ks2k[3][3];
69+
pair_alterpol(scrtyp, r, scaleb, cut, off, xr, yr, zr, @springi@, @sizi@, @alphai@,
70+
springk, sizk, alphak, ks2i, ks2k);
71+
@psci00@ += ks2i[0][0];
72+
@psci01@ += ks2i[0][1];
73+
@psci02@ += ks2i[0][2];
74+
@psci10@ += ks2i[1][0];
75+
@psci11@ += ks2i[1][1];
76+
@psci12@ += ks2i[1][2];
77+
@psci20@ += ks2i[2][0];
78+
@psci21@ += ks2i[2][1];
79+
@psci22@ += ks2i[2][2];
80+
psck00 += ks2k[0][0];
81+
psck01 += ks2k[0][1];
82+
psck02 += ks2k[0][2];
83+
psck10 += ks2k[1][0];
84+
psck11 += ks2k[1][1];
85+
psck12 += ks2k[1][2];
86+
psck20 += ks2k[2][0];
87+
psck21 += ks2k[2][1];
88+
psck22 += ks2k[2][2];
89+
}

ext/ext/yaml/dexpol.yaml

Lines changed: 100 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,100 @@
1+
KERNEL_NAME: dexpol_cu1
2+
TEMPLATE_PARAMS: template<class Ver>
3+
CONSTEXPR_FLAGS: |
4+
constexpr bool do_v = Ver::v;
5+
6+
VIRIAL:
7+
- vep
8+
GRADIENT:
9+
- gx
10+
- gy
11+
- gz
12+
13+
CUT_DISTANCE:
14+
- 'cut'
15+
16+
OFF_DISTANCE:
17+
- 'off'
18+
19+
EXCLUDE_INFO:
20+
- 'dinfo'
21+
22+
SCALE_1X_TYPE: real3_const_array,1
23+
24+
EXTRA_PARAMS: |
25+
, const real* restrict polarity, const real (*restrict uind)[3]
26+
, const real* restrict kpep, const real* restrict prepep
27+
, const real* restrict dmppep, const int* restrict lpep
28+
, ExpolScr scrtyp, real f
29+
30+
I_POSITION:
31+
- def: shared real xi from:x
32+
- def: shared real yi from:y
33+
- def: shared real zi from:z
34+
K_POSITION:
35+
- def: register real xk from:x
36+
- def: register real yk from:y
37+
- def: register real zk from:z
38+
39+
I_FORCE:
40+
- def: shared real frcxi addto:gx
41+
- def: shared real frcyi addto:gy
42+
- def: shared real frczi addto:gz
43+
44+
K_FORCE:
45+
- def: register real frcxk addto:gx
46+
- def: register real frcyk addto:gy
47+
- def: register real frczk addto:gz
48+
49+
I_VARIABLES:
50+
- def: shared real uix from:uind,0
51+
- def: shared real uiy from:uind,1
52+
- def: shared real uiz from:uind,2
53+
- def: shared real springi from:kpep
54+
- def: shared real sizi from:prepep
55+
- def: shared real alphai from:dmppep
56+
- def: shared int epli from:lpep
57+
- def: shared real poli from:polarity
58+
59+
K_VARIABLES:
60+
- def: register real ukx from:uind,0
61+
- def: register real uky from:uind,1
62+
- def: register real ukz from:uind,2
63+
- def: register real springk from:kpep
64+
- def: register real sizk from:prepep
65+
- def: register real alphak from:dmppep
66+
- def: register int eplk from:lpep
67+
- def: register real polk from:polarity
68+
69+
FULL_PAIRWISE_INTERACTION: |
70+
real xr = xk - @xi@;
71+
real yr = yk - @yi@;
72+
real zr = zk - @zi@;
73+
real r2 = image2(xr, yr, zr);
74+
if ((eplk or @epli@) and r2 <= off * off and incl) {
75+
real r = REAL_SQRT(r2);
76+
real frc[3];
77+
pair_dexpol(scrtyp, r, scaleb, cut, off, xr, yr, zr, @uix@, @uiy@, @uiz@, ukx, uky, ukz,
78+
@springi@/@poli@, @sizi@, @alphai@, springk/polk, sizk, alphak, f, frc);
79+
@frcxi@ += frc[0];
80+
@frcyi@ += frc[1];
81+
@frczi@ += frc[2];
82+
frcxk -= frc[0];
83+
frcyk -= frc[1];
84+
frczk -= frc[2];
85+
86+
if CONSTEXPR (do_v) {
87+
real vxx = -xr * frc[0];
88+
real vxy = -0.5f * (yr * frc[0] + xr * frc[1]);
89+
real vxz = -0.5f * (zr * frc[0] + xr * frc[2]);
90+
real vyy = -yr * frc[1];
91+
real vyz = -0.5f * (zr * frc[1] + yr * frc[2]);
92+
real vzz = -zr * frc[2];
93+
veptlxx += floatTo<vbuf_prec>(vxx);
94+
veptlyx += floatTo<vbuf_prec>(vxy);
95+
veptlzx += floatTo<vbuf_prec>(vxz);
96+
veptlyy += floatTo<vbuf_prec>(vyy);
97+
veptlzy += floatTo<vbuf_prec>(vyz);
98+
veptlzz += floatTo<vbuf_prec>(vzz);
99+
}
100+
}

include/ff/cuinduce.h

Lines changed: 30 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,30 @@
1+
#pragma once
2+
#include "ff/precision.h"
3+
4+
namespace tinker {
5+
// udir = polarity * field
6+
7+
__global__
8+
void pcgUdirV1(int n, const real* polarity, //
9+
real (*udir)[3], const real (*field)[3]);
10+
11+
__global__
12+
void pcgUdirV2(int n, const real* polarity, //
13+
real (*udir)[3], real (*udirp)[3], const real (*field)[3], const real (*fieldp)[3]);
14+
15+
// r(0) = E - (1/polarity + Tu) u(0) = (udir - u(0))/polarity + mutual field
16+
17+
__global__
18+
void pcgRsd0V1(int n, const real* polarity_inv, real (*rsd)[3], //
19+
const real (*udir)[3], const real (*uind)[3], const real (*field)[3]);
20+
21+
__global__
22+
void pcgRsd0V2(int n, const real* polarity_inv, real (*rsd)[3], real (*rsp)[3], //
23+
const real (*udir)[3], const real (*udip)[3], const real (*uind)[3], const real (*uinp)[3],
24+
const real (*field)[3], const real (*fielp)[3]);
25+
26+
__global__
27+
void pcgRsd0V3(int n, const real* polarity_inv, real (*rsd)[3], //
28+
const real (*udir)[3], const real (*uind)[3], const real (*field)[3],
29+
const real (*polscale)[3][3]);
30+
}

include/ff/hippo/expol.h

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
#pragma once
2+
#include "ff/energybuffer.h"
3+
#include "tool/rcman.h"
4+
5+
namespace tinker {
6+
void expolData(RcOp);
7+
8+
void alterpol(real (*polscale)[3][3], real (*polinv)[3][3]);
9+
void dexpol(int vers);
10+
}

include/ff/hippo/expolscr.h

Lines changed: 11 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,11 @@
1+
#pragma once
2+
3+
namespace tinker {
4+
enum class ExpolScr
5+
{
6+
NONE,
7+
S2U,
8+
S2,
9+
G,
10+
};
11+
}

include/ff/hippomod.h

Lines changed: 12 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -2,6 +2,7 @@
22
#include "ff/energybuffer.h"
33
#include "ff/hippo/chgpen.h"
44
#include "ff/hippo/echgtrn.h"
5+
#include "ff/hippo/expolscr.h"
56

67
// mplpot
78
namespace tinker {
@@ -40,10 +41,6 @@ TINKER_EXTERN int nmdwexclude;
4041
TINKER_EXTERN int (*mdwexclude)[2];
4142
TINKER_EXTERN real (*mdwexclude_scale)[3];
4243

43-
TINKER_EXTERN int ndexclude;
44-
TINKER_EXTERN int (*dexclude)[2];
45-
TINKER_EXTERN real* dexclude_scale;
46-
4744
TINKER_EXTERN int nwexclude;
4845
TINKER_EXTERN int (*wexclude)[2];
4946
TINKER_EXTERN real* wexclude_scale;
@@ -65,3 +62,14 @@ TINKER_EXTERN virial_prec virial_ect[9];
6562

6663
TINKER_EXTERN Chgtrn ctrntyp;
6764
}
65+
66+
// expol
67+
namespace tinker {
68+
TINKER_EXTERN real* kpep;
69+
TINKER_EXTERN real* prepep;
70+
TINKER_EXTERN real* dmppep;
71+
TINKER_EXTERN int* lpep;
72+
TINKER_EXTERN real (*polscale)[3][3];
73+
TINKER_EXTERN real (*polinv)[3][3];
74+
TINKER_EXTERN ExpolScr scrtyp;
75+
}

0 commit comments

Comments
 (0)