kCalculateCDLJEwaldFastReciprocal.h 6.17 KB
Newer Older
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
/* -------------------------------------------------------------------------- *
 *                                   OpenMM                                   *
 * -------------------------------------------------------------------------- *
 * This is part of the OpenMM molecular simulation toolkit originating from   *
 * Simbios, the NIH National Center for Physics-Based Simulation of           *
 * Biological Structures at Stanford, funded under the NIH Roadmap for        *
 * Medical Research, grant U54 GM072970. See https://simtk.org.               *
 *                                                                            *
 * Portions copyright (c) 2009 Stanford University and the Authors.           *
 * Authors: Rossen P. Apostolov, Peter Eastman                                     *
 * Contributors:                                                              *
 *                                                                            *
 * This program is free software: you can redistribute it and/or modify       *
 * it under the terms of the GNU Lesser General Public License as published   *
 * by the Free Software Foundation, either version 3 of the License, or       *
 * (at your option) any later version.                                        *
 *                                                                            *
 * This program is distributed in the hope that it will be useful,            *
 * but WITHOUT ANY WARRANTY; without even the implied warranty of             *
 * MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.  See the              *
 * GNU Lesser General Public License for more details.                        *
 *                                                                            *
 * You should have received a copy of the GNU Lesser General Public License   *
 * along with this program.  If not, see <http://www.gnu.org/licenses/>.      *
 * -------------------------------------------------------------------------- */

27
28
29
30
31
/**
 * This file contains the kernel for evaluating nonbonded forces using the
 * Ewald summation method (Reciprocal space summation).
 */

32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
/* Define multiply operations for floats */

__device__ float2 MultofFloat2(float2 a, float2 b)
{
    float2 c;
    c.x = a.x * b.x - a.y * b.y;
    c.y = a.x * b.y + a.y * b.x;
    return c;
}

__device__ float2 ConjMultofFloat2(float2 a, float2 b)
{
    float2 c;
    c.x = a.x*b.x + a.y*b.y;
    c.y = a.y*b.x - a.x*b.y;
    return c;
}

50
51
52
/**
 * Precompute the cosine and sine sums which appear in each force term.
 */
53

54
55
__global__ void kCalculateEwaldFastCosSinSums_kernel()
{
56
57
58
59
    const unsigned int ksizex = 2*cSim.kmaxX-1;
    const unsigned int ksizey = 2*cSim.kmaxY-1;
    const unsigned int ksizez = 2*cSim.kmaxZ-1;
    const unsigned int totalK = ksizex*ksizey*ksizez;
60
61
    unsigned int index = threadIdx.x + blockIdx.x * blockDim.x;
    while (index < totalK)
62
    {
63
64
        // Find the wave vector (kx, ky, kz) this index corresponds to.

65
66
67
68
69
        int rx = index/(ksizey*ksizez);
        int remainder = index - rx*ksizey*ksizez;
        int ry = remainder/ksizez;
        int rz = remainder - ry*ksizez - cSim.kmaxZ + 1;
        ry += -cSim.kmaxY + 1;
70
71
72
73
74
75
        float kx = rx*cSim.recipBoxSizeX;
        float ky = ry*cSim.recipBoxSizeY;
        float kz = rz*cSim.recipBoxSizeZ;

        // Compute the sum for this wave vector.

76
77
78
        float2 sum = make_float2(0.0f, 0.0f);
        for (int atom = 0; atom < cSim.atoms; atom++)
        {
79
80
81
82
83
84
85
86
87
            float4 apos = cSim.pPosq[atom];
            float phase = apos.x*kx;
            float2 structureFactor = make_float2(cos(phase), sin(phase));
            phase = apos.y*ky;
            structureFactor = MultofFloat2(structureFactor, make_float2(cos(phase), sin(phase)));
            phase = apos.z*kz;
            structureFactor = MultofFloat2(structureFactor, make_float2(cos(phase), sin(phase)));
            sum.x += apos.w*structureFactor.x;
            sum.y += apos.w*structureFactor.y;
88
        }
89
90
        cSim.pEwaldCosSinSum[index] = sum;
        index += blockDim.x * gridDim.x;
91
92
93
    }
}

94
95
96
97
98
/**
 * Compute the reciprocal space part of the Ewald force, using the precomputed sums from the
 * previous routine.
 */

99
100
101
__global__ void kCalculateEwaldFastForces_kernel()
{

102
103
104
    float PI = 3.14159265358979323846f;
    const float epsilon =  1.0;
    float recipCoeff = cSim.epsfac*(4*PI/cSim.cellVolume/epsilon);
105

106
    unsigned int atom = threadIdx.x + blockIdx.x * blockDim.x;
107
108
109

    while (atom < cSim.atoms)
    {
110
        float4 force = cSim.pForce4[atom];
111
112
113
114
115
116
        float4 apos = cSim.pPosq[atom];

        // Loop over all wave vectors.

        int lowry = 0;
        int lowrz = 1;
117
        for (int rx = 0; rx < cSim.kmaxX; rx++) {
118
            float kx = rx * cSim.recipBoxSizeX;
119
            for (int ry = lowry; ry < cSim.kmaxY; ry++) {
120
                float ky = ry * cSim.recipBoxSizeY;
121
122
123
124
                float phase = apos.x*kx;
                float2 tab_xy = make_float2(cos(phase), sin(phase));
                phase = apos.y*ky;
                tab_xy = MultofFloat2(tab_xy, make_float2(cos(phase), sin(phase)));
125
                for (int rz = lowrz; rz < cSim.kmaxZ; rz++) {
126
                    float kz = rz * cSim.recipBoxSizeZ;
127
128
129

                    // Compute the force contribution of this wave vector.

130
                    int index = rx*(cSim.kmaxY*2-1)*(cSim.kmaxZ*2-1) + (ry+cSim.kmaxY-1)*(cSim.kmaxZ*2-1) + (rz+cSim.kmaxZ-1);
131
132
                    float k2 = kx*kx + ky*ky + kz*kz;
                    float ak = exp(k2*cSim.factorEwald) / k2;
133
134
                    phase = apos.z*kz;
                    float2 structureFactor = MultofFloat2(tab_xy, make_float2(cos(phase), sin(phase)));
135
                    float2 cosSinSum = cSim.pEwaldCosSinSum[index];
136
                    float dEdR = ak * apos.w * (cosSinSum.x*structureFactor.y - cosSinSum.y*structureFactor.x);
137
138
139
                    force.x += 2 * recipCoeff * dEdR * kx;
                    force.y += 2 * recipCoeff * dEdR * ky;
                    force.z += 2 * recipCoeff * dEdR * kz;
140
                    lowrz = 1 - cSim.kmaxZ;
141
                }
142
                lowry = 1 - cSim.kmaxY;
143
            }
144
        }
145
146
147

        // Record the force on the atom.

148
        cSim.pForce4[atom] = force;
149
        atom += blockDim.x * gridDim.x;
150
151
    }
}