kCalculateAmoebaCudaPmeMutualInducedField.h 10.2 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
27
28
29
30
31
/* -------------------------------------------------------------------------- *
 *                                   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: Scott Le Grand, 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/>.      *
 * -------------------------------------------------------------------------- */

#include "amoebaScaleFactors.h"

__global__
#if (__CUDA_ARCH__ >= 200)
__launch_bounds__(GF1XX_NONBOND_THREADS_PER_BLOCK, 1)
Peter Eastman's avatar
Peter Eastman committed
32
#elif (__CUDA_ARCH__ >= 120)
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
__launch_bounds__(GT2XX_NONBOND_THREADS_PER_BLOCK, 1)
#else
__launch_bounds__(G8X_NONBOND_THREADS_PER_BLOCK, 1)
#endif
void METHOD_NAME(kCalculateAmoebaPmeMutualInducedField, _kernel)(
                            unsigned int* workUnit,
                            float* outputField, float* outputFieldPolar
){

    extern __shared__ MutualInducedParticle sA[];

    unsigned int totalWarps      = gridDim.x*blockDim.x/GRID;
    unsigned int warp            = (blockIdx.x*blockDim.x+threadIdx.x)/GRID;
    unsigned int numWorkUnits    = cSim.pInteractionCount[0];
    unsigned int pos             = warp*numWorkUnits/totalWarps;
    unsigned int end             = (warp+1)*numWorkUnits/totalWarps;
    unsigned int lasty           = 0xFFFFFFFF;
    const float uscale           = 1.0f;

    while (pos < end)
    {

        unsigned int x;
        unsigned int y;
        bool bExclusionFlag;

        // Extract cell coordinates

        decodeCell( workUnit[pos], &x, &y, &bExclusionFlag );

        unsigned int tgx                 = threadIdx.x & (GRID - 1);
        unsigned int tbx                 = threadIdx.x - tgx;
        unsigned int tj                  = tgx;

        MutualInducedParticle*  psA      = &sA[tbx];
        unsigned int atomI               = x + tgx;
        MutualInducedParticle localParticle;
        loadMutualInducedShared( &localParticle, atomI );

        float fieldSum[3];
        float fieldPolarSum[3];

        // 0: field at i due to j
        // 1: field at i due to j polar

        fieldSum[0]                      = 0.0f;
        fieldSum[1]                      = 0.0f;
        fieldSum[2]                      = 0.0f;

        fieldPolarSum[0]                 = 0.0f;
        fieldPolarSum[1]                 = 0.0f;
        fieldPolarSum[2]                 = 0.0f;

86
        if (x == y ){
87
88
89
90
91

            // load shared data

            loadMutualInducedShared( &(sA[threadIdx.x]), atomI );

92
            for (unsigned int j = 0; j < GRID; j++) {
Mark Friedrichs's avatar
Mark Friedrichs committed
93
                if(  ( (atomI != (y + j)) && (atomI < cSim.atoms) && ((y+j) < cSim.atoms) ) ){
94
95
                    float4 delta;
                    float prefactor2;
Mark Friedrichs's avatar
Mark Friedrichs committed
96
                    setupMutualInducedFieldPairIxn_kernel( localParticle, psA[j], uscale, &delta, &prefactor2 );
Mark Friedrichs's avatar
Mark Friedrichs committed
97
98
                    calculateMutualInducedFieldPairIxn_kernel(  psA[j].inducedDipole,      delta, prefactor2, fieldSum );
                    calculateMutualInducedFieldPairIxn_kernel(  psA[j].inducedDipolePolar, delta, prefactor2, fieldPolarSum );
Mark Friedrichs's avatar
Mark Friedrichs committed
99
                }
Mark Friedrichs's avatar
Mark Friedrichs committed
100

101
102
103
104
105
            }

            // Write results

#ifdef USE_OUTPUT_BUFFER_PER_WARP
106
            unsigned int offset            = 3*(x + tgx + warp*cSim.paddedNumberOfAtoms);
107
#else
108
            unsigned int offset            = 3*(x + tgx + (x >> GRIDBITS) * cSim.paddedNumberOfAtoms);
Mark Friedrichs's avatar
Mark Friedrichs committed
109
#endif
110
111
112
            load3dArray( offset, fieldSum,      outputField );
            load3dArray( offset, fieldPolarSum, outputFieldPolar);

Mark Friedrichs's avatar
Mark Friedrichs committed
113
        } else {
114

115
            if( lasty != y ){
Mark Friedrichs's avatar
Mark Friedrichs committed
116
117
118
119
                unsigned int atomJ        = y + tgx;
                loadMutualInducedShared( &(sA[threadIdx.x]), atomJ );
            }
    
Mark Friedrichs's avatar
Mark Friedrichs committed
120
            unsigned int flags = cSim.pInteractionFlag[pos];
121
            if( flags != 0 ){
122

Mark Friedrichs's avatar
Mark Friedrichs committed
123
124
125
126
#ifndef INCLUDE_MI_FIELD_BUFFERS
                flags = 0xFFFFFFFF;
#endif

Mark Friedrichs's avatar
Mark Friedrichs committed
127
128
129
130
               // zero shared fields
    
                zeroMutualInducedParticleSharedField(  &(sA[threadIdx.x]) );
    
Mark Friedrichs's avatar
Mark Friedrichs committed
131
132
                for (unsigned int j = 0; j < GRID; j++){
                    if ((flags&(1<<j)) != 0) {
Peter Eastman's avatar
Peter Eastman committed
133
                        unsigned int jIdx = (flags == 0xFFFFFFFF) ? tj : j;
Mark Friedrichs's avatar
Mark Friedrichs committed
134
135
                        float4 delta;
                        float prefactor2;
Mark Friedrichs's avatar
Mark Friedrichs committed
136
                        if( (atomI < cSim.atoms) && ((y+jIdx) < cSim.atoms) ){
Mark Friedrichs's avatar
Mark Friedrichs committed
137
                            setupMutualInducedFieldPairIxn_kernel( localParticle, psA[jIdx], uscale, &delta, &prefactor2 );
Mark Friedrichs's avatar
Mark Friedrichs committed
138
139
                            calculateMutualInducedFieldPairIxn_kernel(  psA[jIdx].inducedDipole,          delta, prefactor2, fieldSum );
                            calculateMutualInducedFieldPairIxn_kernel(  psA[jIdx].inducedDipolePolar,     delta, prefactor2, fieldPolarSum );
Mark Friedrichs's avatar
Mark Friedrichs committed
140
141
142
#ifndef INCLUDE_MI_FIELD_BUFFERS
                            calculateMutualInducedFieldPairIxn_kernel(  localParticle.inducedDipole,      delta, prefactor2, psA[jIdx].field );
                            calculateMutualInducedFieldPairIxn_kernel(  localParticle.inducedDipolePolar, delta, prefactor2, psA[jIdx].fieldPolar );
Mark Friedrichs's avatar
Mark Friedrichs committed
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
#else
                            if( flags == 0xFFFFFFFF ){
                                calculateMutualInducedFieldPairIxn_kernel(  localParticle.inducedDipole,      delta, prefactor2, psA[jIdx].field );
                                calculateMutualInducedFieldPairIxn_kernel(  localParticle.inducedDipolePolar, delta, prefactor2, psA[jIdx].fieldPolar );
                            } else {
                                calculateMutualInducedFieldPairIxnNoAdd_kernel(  localParticle.inducedDipole,      delta, prefactor2,  sA[threadIdx.x].tempBuffer );
                                calculateMutualInducedFieldPairIxnNoAdd_kernel(  localParticle.inducedDipolePolar, delta, prefactor2,  sA[threadIdx.x].tempBufferP );
    
                                if( tgx % 2 == 0 ){
                                    sumTempBuffer( sA[threadIdx.x], sA[threadIdx.x+1] );
                                }
                                if( tgx % 4 == 0 ){
                                    sumTempBuffer( sA[threadIdx.x], sA[threadIdx.x+2] );
                                }
                                if( tgx % 8 == 0 ){
                                    sumTempBuffer( sA[threadIdx.x], sA[threadIdx.x+4] );
                                }
                                if( tgx % 16 == 0 ){
                                    sumTempBuffer( sA[threadIdx.x], sA[threadIdx.x+8] );
                                }
    
                                if (tgx == 0)
                                {
                                    psA[jIdx].field[0]         += sA[threadIdx.x].tempBuffer[0]  + sA[threadIdx.x+16].tempBuffer[0];
                                    psA[jIdx].field[1]         += sA[threadIdx.x].tempBuffer[1]  + sA[threadIdx.x+16].tempBuffer[1];
                                    psA[jIdx].field[2]         += sA[threadIdx.x].tempBuffer[2]  + sA[threadIdx.x+16].tempBuffer[2];
    
                                    psA[jIdx].fieldPolar[0]    += sA[threadIdx.x].tempBufferP[0] + sA[threadIdx.x+16].tempBufferP[0];
                                    psA[jIdx].fieldPolar[1]    += sA[threadIdx.x].tempBufferP[1] + sA[threadIdx.x+16].tempBufferP[1];
                                    psA[jIdx].fieldPolar[2]    += sA[threadIdx.x].tempBufferP[2] + sA[threadIdx.x+16].tempBufferP[2];
                                }
    
Peter Eastman's avatar
Peter Eastman committed
175
                            }
176
#endif
Mark Friedrichs's avatar
Mark Friedrichs committed
177
                        }
Peter Eastman's avatar
Peter Eastman committed
178
                    }
Mark Friedrichs's avatar
Mark Friedrichs committed
179
    
180
                    tj = (tj + 1) & (GRID - 1);
Mark Friedrichs's avatar
Mark Friedrichs committed
181
182
183
184
185
    
                } // end of j-loop
    
                // Write results
    
186
#ifdef USE_OUTPUT_BUFFER_PER_WARP
187
                unsigned int offset     = 3*(x + tgx + warp*cSim.paddedNumberOfAtoms);
Mark Friedrichs's avatar
Mark Friedrichs committed
188
189
190
                load3dArrayBufferPerWarp( offset, fieldSum,      outputField );
                load3dArrayBufferPerWarp( offset, fieldPolarSum, outputFieldPolar);
    
191
                offset                  = 3*(y + tgx + warp*cSim.paddedNumberOfAtoms);
Mark Friedrichs's avatar
Mark Friedrichs committed
192
193
194
195
    
                load3dArrayBufferPerWarp( offset, sA[threadIdx.x].field,      outputField );
                load3dArrayBufferPerWarp( offset, sA[threadIdx.x].fieldPolar, outputFieldPolar);
    
196
#else
197
                unsigned int offset     = 3*(x + tgx + (y >> GRIDBITS) * cSim.paddedNumberOfAtoms);
Mark Friedrichs's avatar
Mark Friedrichs committed
198
199
200
                load3dArray( offset, fieldSum,      outputField );
                load3dArray( offset, fieldPolarSum, outputFieldPolar);
    
201
                offset                  = 3*(y + tgx + (x >> GRIDBITS) * cSim.paddedNumberOfAtoms);
Mark Friedrichs's avatar
Mark Friedrichs committed
202
203
204
                load3dArray( offset, sA[threadIdx.x].field,      outputField );
                load3dArray( offset, sA[threadIdx.x].fieldPolar, outputFieldPolar);
    
205
#endif
Mark Friedrichs's avatar
Mark Friedrichs committed
206
207
208
                lasty = y;
    
            } // end of pInteractionFlag block
209

Mark Friedrichs's avatar
Mark Friedrichs committed
210
        } // end of x == y block
211
212
213
        pos++;
    }
}