Skip to content

GitLab

Commit 50a15fb0 authored by Mark Friedrichs's avatar Mark Friedrichs
Browse files

Mods to reduce calculation of ixns within blocks for Vdw

parent e6c19b54
Hide whitespace changes
Inline Side-by-side
Showing with 121 additions and 71 deletions
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaPmeDirectElectrostatic.h
View file @ 50a15fb0
...@@ -268,14 +268,6 @@ if( atomI == targetAtom ){ ...@@ -268,14 +268,6 @@ if( atomI == targetAtom ){
// No interactions in this block. // No interactions in this block.
} else { } else {
if (lasty != y) {
// load shared data
loadPmeDirectElectrostaticShared( &(sA[threadIdx.x]), (y+tgx) );
}
sA[threadIdx.x].force[0] = 0.0f; sA[threadIdx.x].force[0] = 0.0f;
sA[threadIdx.x].force[1] = 0.0f; sA[threadIdx.x].force[1] = 0.0f;
sA[threadIdx.x].force[2] = 0.0f; sA[threadIdx.x].force[2] = 0.0f;
......
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaVdw14_7.cu
View file @ 50a15fb0
...@@ -250,6 +250,12 @@ void kCalculateAmoebaVdw14_7Reduction_kernel( float* inputForce, float4* outputF ...@@ -250,6 +250,12 @@ void kCalculateAmoebaVdw14_7Reduction_kernel( float* inputForce, float4* outputF
} }
} }
__device__ void sumTempBuffer( Vdw14_7Particle& atomI, Vdw14_7Particle& atomJ ){
atomI.tempForce[0] += atomJ.tempForce[0];
atomI.tempForce[1] += atomJ.tempForce[1];
atomI.tempForce[2] += atomJ.tempForce[2];
}
static void kCalculateAmoebaVdw14_7Reduction(amoebaGpuContext amoebaGpu, CUDAStream<float>* vdwOutputArray, CUDAStream<float4>* forceOutputArray ) static void kCalculateAmoebaVdw14_7Reduction(amoebaGpuContext amoebaGpu, CUDAStream<float>* vdwOutputArray, CUDAStream<float4>* forceOutputArray )
{ {
......
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaVdw14_7.h
View file @ 50a15fb0
...@@ -212,79 +212,126 @@ if( atomI == targetAtom || (y+j) == targetAtom ){ ...@@ -212,79 +212,126 @@ if( atomI == targetAtom || (y+j) == targetAtom ){
if (flags == 0) { if (flags == 0) {
} else { } else {
#endif #endif
// zero shared fields
// zero shared fields
zeroVdw14_7SharedForce( &(sA[threadIdx.x]) );
if( bExclusionFlag ){
unsigned int xi = x >> GRIDBITS;
unsigned int yi = y >> GRIDBITS;
unsigned int cell = xi+yi*cSim.paddedNumberOfAtoms/GRID-yi*(yi+1)/2;
exclusionIndex = cAmoebaSim.pVdwExclusionIndicesIndex[cell]+tgx;
exclusionMask = cAmoebaSim.pVdwExclusionIndices[exclusionIndex];
}
for (unsigned int j = 0; j < GRID; j++)
{
float ijForce[3]; zeroVdw14_7SharedForce( &(sA[threadIdx.x]) );
// get combined sigma and epsilon if( bExclusionFlag )
{
float combindedSigma; unsigned int xi = x >> GRIDBITS;
float combindedEpsilon; unsigned int yi = y >> GRIDBITS;
getVdw14_7CombindedSigmaEpsilon_kernel( sigmaCombiningRule, localParticle.sigma, psA[tj].sigma, &combindedSigma, unsigned int cell = xi+yi*cSim.paddedNumberOfAtoms/GRID-yi*(yi+1)/2;
epsilonCombiningRule, localParticle.epsilon, psA[tj].epsilon, &combindedEpsilon );
// calculate force exclusionIndex = cAmoebaSim.pVdwExclusionIndicesIndex[cell]+tgx;
exclusionMask = cAmoebaSim.pVdwExclusionIndices[exclusionIndex];
}
float energy; for (unsigned int j = 0; j < GRID; j++)
ijForce[0] = psA[tj].x - localParticle.x;
ijForce[1] = psA[tj].y - localParticle.y;
ijForce[2] = psA[tj].z - localParticle.z;
if( cAmoebaSim.vdwUsePBC )
{ {
ijForce[0] -= floor(ijForce[0]*cSim.invPeriodicBoxSizeX+0.5f)*cSim.periodicBoxSizeX;
ijForce[1] -= floor(ijForce[1]*cSim.invPeriodicBoxSizeY+0.5f)*cSim.periodicBoxSizeY; float ijForce[3];
ijForce[2] -= floor(ijForce[2]*cSim.invPeriodicBoxSizeZ+0.5f)*cSim.periodicBoxSizeZ; #ifdef USE_CUTOFF
} if ((flags&(1<<j)) != 0)
calculateVdw14_7PairIxn_kernel( combindedSigma, combindedEpsilon, ijForce, &energy {
#ifdef AMOEBA_DEBUG unsigned int jIdx = (flags == 0xFFFFFFFF) ? tj : j;
, pullDebug #else
unsigned int jIdx = tj;
#endif #endif
);
// get combined sigma and epsilon
// mask out excluded ixns
unsigned int mask = ( (atomI >= cAmoebaSim.numberOfAtoms) || ((y+tj) >= cAmoebaSim.numberOfAtoms) ) ? 0 : 1; float combindedSigma;
if( mask && bExclusionFlag ){ float combindedEpsilon;
unsigned int maskIndex = 1 << tj; getVdw14_7CombindedSigmaEpsilon_kernel( sigmaCombiningRule, localParticle.sigma, psA[jIdx].sigma, &combindedSigma,
mask = (exclusionMask & maskIndex) ? 0 : 1; epsilonCombiningRule, localParticle.epsilon, psA[jIdx].epsilon, &combindedEpsilon );
}
// calculate force
// accumulate force for atomI
forceSum[0] += mask ? ijForce[0] : 0.0f;
forceSum[1] += mask ? ijForce[1] : 0.0f;
forceSum[2] += mask ? ijForce[2] : 0.0f;
// accumulate force for atomJ float energy;
ijForce[0] = psA[jIdx].x - localParticle.x;
ijForce[1] = psA[jIdx].y - localParticle.y;
ijForce[2] = psA[jIdx].z - localParticle.z;
if( cAmoebaSim.vdwUsePBC )
{
ijForce[0] -= floor(ijForce[0]*cSim.invPeriodicBoxSizeX+0.5f)*cSim.periodicBoxSizeX;
ijForce[1] -= floor(ijForce[1]*cSim.invPeriodicBoxSizeY+0.5f)*cSim.periodicBoxSizeY;
ijForce[2] -= floor(ijForce[2]*cSim.invPeriodicBoxSizeZ+0.5f)*cSim.periodicBoxSizeZ;
}
calculateVdw14_7PairIxn_kernel( combindedSigma, combindedEpsilon, ijForce, &energy
#ifdef AMOEBA_DEBUG
, pullDebug
#endif
);
// mask out excluded ixns
psA[tj].force[0] -= mask ? ijForce[0] : 0.0f; unsigned int mask = ( (atomI >= cAmoebaSim.numberOfAtoms) || ((y+jIdx) >= cAmoebaSim.numberOfAtoms) ) ? 0 : 1;
psA[tj].force[1] -= mask ? ijForce[1] : 0.0f; if( mask && bExclusionFlag ){
psA[tj].force[2] -= mask ? ijForce[2] : 0.0f; unsigned int maskIndex = 1 << jIdx;
mask = (exclusionMask & maskIndex) ? 0 : 1;
totalEnergy += mask ? energy : 0.0f; }
if( mask == 0 )
{
energy = ijForce[0] = ijForce[1] = ijForce[2] = 0.0f;
}
// accumulate force for atomI
forceSum[0] += ijForce[0];
forceSum[1] += ijForce[1];
forceSum[2] += ijForce[2];
totalEnergy += energy;
#ifndef USE_CUTOFF
psA[jIdx].force[0] -= ijForce[0];
psA[jIdx].force[1] -= ijForce[1];
psA[jIdx].force[2] -= ijForce[2];
#else
if( flags == 0xFFFFFFFF ){
psA[jIdx].force[0] -= ijForce[0];
psA[jIdx].force[1] -= ijForce[1];
psA[jIdx].force[2] -= ijForce[2];
} else {
sA[threadIdx.x].tempForce[0] = ijForce[0];
sA[threadIdx.x].tempForce[1] = ijForce[1];
sA[threadIdx.x].tempForce[2] = ijForce[2];
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].force[0] -= sA[threadIdx.x].tempForce[0] + sA[threadIdx.x+16].tempForce[0];
psA[jIdx].force[1] -= sA[threadIdx.x].tempForce[1] + sA[threadIdx.x+16].tempForce[1];
psA[jIdx].force[2] -= sA[threadIdx.x].tempForce[2] + sA[threadIdx.x+16].tempForce[2];
}
}
#endif
#ifdef AMOEBA_DEBUG #ifdef AMOEBA_DEBUG
if( atomI == targetAtom || (y+tj) == targetAtom ){ if( atomI == targetAtom || (y+jIdx) == targetAtom ){
unsigned int index = (atomI == targetAtom) ? (y + tj) : atomI; unsigned int index = (atomI == targetAtom) ? (y + jIdx) : atomI;
debugArray[index].x = (float) atomI; debugArray[index].x = (float) atomI;
debugArray[index].y = (float) (y + tj); debugArray[index].y = (float) (y + jIdx);
debugArray[index].z = -3.0; debugArray[index].z = -3.0;
debugArray[index].w = (float) (mask + 1); debugArray[index].w = (float) (mask + 1);
...@@ -312,11 +359,15 @@ if( atomI == targetAtom || (y+tj) == targetAtom ){ ...@@ -312,11 +359,15 @@ if( atomI == targetAtom || (y+tj) == targetAtom ){
debugArray[index].z = mask ? ijForce[2] : 0.0f; debugArray[index].z = mask ? ijForce[2] : 0.0f;
} }
#endif #endif
tj = (tj + 1) & (GRID - 1);
} // end of j-loop
#ifdef USE_CUTOFF #ifdef USE_CUTOFF
} }
#endif
tj = (tj + 1) & (GRID - 1);
} // end of j-loop
#ifdef USE_CUTOFF
}
#endif #endif
// Write results // Write results
......
plugins/amoeba/platforms/cuda/src/kernels/kCalculateAmoebaCudaVdwParticle.h
View file @ 50a15fb0
...@@ -13,6 +13,7 @@ struct Vdw14_7Particle { ...@@ -13,6 +13,7 @@ struct Vdw14_7Particle {
float epsilon; float epsilon;
float force[3]; float force[3];
float tempForce[3];
}; };
......
Markdown is supported
0% or .
You are about to add 0 people to the discussion. Proceed with caution.
Finish editing this message first!
Please register or to comment