Skip to content

GitLab

Commit 4d51fb04 authored by peastman's avatar peastman
Browse files

Merge pull request #407 from peastman/master 

Fixed bug in OpenCL implementation of GBSA on CPU devices
parents 71d91890 cd0ee42e
Show whitespace changes
Inline Side-by-side
Showing with 21 additions and 4 deletions
platforms/opencl/src/kernels/gbsaObc_cpu.cl
View file @ 4d51fb04
...@@ -468,11 +468,16 @@ __kernel void computeGBSAForce1( ...@@ -468,11 +468,16 @@ __kernel void computeGBSAForce1(
real expTerm = EXP(-D_ij); real expTerm = EXP(-D_ij);
real denominator2 = r2 + alpha2_ij*expTerm; real denominator2 = r2 + alpha2_ij*expTerm;
real denominator = SQRT(denominator2); real denominator = SQRT(denominator2);
real tempEnergy = (PREFACTOR*posq1.w*posq2.w)*RECIP(denominator); real scaledChargeProduct = PREFACTOR*posq1.w*posq2.w;
real tempEnergy = scaledChargeProduct*RECIP(denominator);
real Gpol = tempEnergy*RECIP(denominator2); real Gpol = tempEnergy*RECIP(denominator2);
real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij); real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij);
real dEdR = Gpol*(1.0f - 0.25f*expTerm); real dEdR = Gpol*(1.0f - 0.25f*expTerm);
force.w += dGpol_dalpha2_ij*bornRadius2; force.w += dGpol_dalpha2_ij*bornRadius2;
#ifdef USE_CUTOFF
if (atom1 != y*TILE_SIZE+j)
tempEnergy -= scaledChargeProduct/CUTOFF;
#endif
energy += 0.5f*tempEnergy; energy += 0.5f*tempEnergy;
delta.xyz *= dEdR; delta.xyz *= dEdR;
force.xyz -= delta.xyz; force.xyz -= delta.xyz;
...@@ -527,11 +532,15 @@ __kernel void computeGBSAForce1( ...@@ -527,11 +532,15 @@ __kernel void computeGBSAForce1(
real expTerm = EXP(-D_ij); real expTerm = EXP(-D_ij);
real denominator2 = r2 + alpha2_ij*expTerm; real denominator2 = r2 + alpha2_ij*expTerm;
real denominator = SQRT(denominator2); real denominator = SQRT(denominator2);
real tempEnergy = (PREFACTOR*posq1.w*posq2.w)*RECIP(denominator); real scaledChargeProduct = PREFACTOR*posq1.w*posq2.w;
real tempEnergy = scaledChargeProduct*RECIP(denominator);
real Gpol = tempEnergy*RECIP(denominator2); real Gpol = tempEnergy*RECIP(denominator2);
real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij); real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij);
real dEdR = Gpol*(1.0f - 0.25f*expTerm); real dEdR = Gpol*(1.0f - 0.25f*expTerm);
force.w += dGpol_dalpha2_ij*bornRadius2; force.w += dGpol_dalpha2_ij*bornRadius2;
#ifdef USE_CUTOFF
tempEnergy -= scaledChargeProduct/CUTOFF;
#endif
energy += tempEnergy; energy += tempEnergy;
delta.xyz *= dEdR; delta.xyz *= dEdR;
force.xyz -= delta.xyz; force.xyz -= delta.xyz;
...@@ -684,11 +693,15 @@ __kernel void computeGBSAForce1( ...@@ -684,11 +693,15 @@ __kernel void computeGBSAForce1(
real expTerm = EXP(-D_ij); real expTerm = EXP(-D_ij);
real denominator2 = r2 + alpha2_ij*expTerm; real denominator2 = r2 + alpha2_ij*expTerm;
real denominator = SQRT(denominator2); real denominator = SQRT(denominator2);
real tempEnergy = (PREFACTOR*posq1.w*posq2.w)*RECIP(denominator); real scaledChargeProduct = PREFACTOR*posq1.w*posq2.w;
real tempEnergy = scaledChargeProduct*RECIP(denominator);
real Gpol = tempEnergy*RECIP(denominator2); real Gpol = tempEnergy*RECIP(denominator2);
real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij); real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij);
real dEdR = Gpol*(1.0f - 0.25f*expTerm); real dEdR = Gpol*(1.0f - 0.25f*expTerm);
force.w += dGpol_dalpha2_ij*bornRadius2; force.w += dGpol_dalpha2_ij*bornRadius2;
#ifdef USE_CUTOFF
tempEnergy -= scaledChargeProduct/CUTOFF;
#endif
energy += tempEnergy; energy += tempEnergy;
delta.xyz *= dEdR; delta.xyz *= dEdR;
force.xyz -= delta.xyz; force.xyz -= delta.xyz;
...@@ -744,11 +757,15 @@ __kernel void computeGBSAForce1( ...@@ -744,11 +757,15 @@ __kernel void computeGBSAForce1(
real expTerm = EXP(-D_ij); real expTerm = EXP(-D_ij);
real denominator2 = r2 + alpha2_ij*expTerm; real denominator2 = r2 + alpha2_ij*expTerm;
real denominator = SQRT(denominator2); real denominator = SQRT(denominator2);
real tempEnergy = (PREFACTOR*posq1.w*posq2.w)*RECIP(denominator); real scaledChargeProduct = PREFACTOR*posq1.w*posq2.w;
real tempEnergy = scaledChargeProduct*RECIP(denominator);
real Gpol = tempEnergy*RECIP(denominator2); real Gpol = tempEnergy*RECIP(denominator2);
real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij); real dGpol_dalpha2_ij = -0.5f*Gpol*expTerm*(1.0f+D_ij);
real dEdR = Gpol*(1.0f - 0.25f*expTerm); real dEdR = Gpol*(1.0f - 0.25f*expTerm);
force.w += dGpol_dalpha2_ij*bornRadius2; force.w += dGpol_dalpha2_ij*bornRadius2;
#ifdef USE_CUTOFF
tempEnergy -= scaledChargeProduct/CUTOFF;
#endif
energy += tempEnergy; energy += tempEnergy;
delta.xyz *= dEdR; delta.xyz *= dEdR;
force.xyz -= delta.xyz; force.xyz -= delta.xyz;
......
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