Skip to content

GitLab

Commit b3dbaeaa authored by Mark Friedrichs's avatar Mark Friedrichs
Browse files

Bug fixes for Nvidia OBC

parent 3a55714e
Show whitespace changes
Inline Side-by-side
Showing with 4 additions and 4 deletions
platforms/opencl/src/kernels/gbsaObc2.cl
View file @ b3dbaeaa
...@@ -17,7 +17,7 @@ ...@@ -17,7 +17,7 @@
float term1 = (0.5f*(0.25f+obcParams2.y*obcParams2.y*invRSquaredOver4)*t2J + t1J*invRSquaredOver4)*invR; float term1 = (0.5f*(0.25f+obcParams2.y*obcParams2.y*invRSquaredOver4)*t2J + t1J*invRSquaredOver4)*invR;
float term2 = (0.5f*(0.25f+obcParams1.y*obcParams1.y*invRSquaredOver4)*t2I + t1I*invRSquaredOver4)*invR; float term2 = (0.5f*(0.25f+obcParams1.y*obcParams1.y*invRSquaredOver4)*t2I + t1I*invRSquaredOver4)*invR;
float tempdEdR = select(0.0f, bornForce1*term1, obcParams1.x < rScaledRadiusJ); float tempdEdR = select(0.0f, bornForce1*term1, obcParams1.x < rScaledRadiusJ);
tempdEdR += select(0.0f, bornForce2*term2, obcParams2.x < rScaledRadiusJ); tempdEdR += select(0.0f, bornForce2*term2, obcParams2.x < rScaledRadiusI);
#ifdef USE_CUTOFF #ifdef USE_CUTOFF
unsigned int includeInteraction = (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS && atom1 != atom2 && r2 < CUTOFF_SQUARED); unsigned int includeInteraction = (atom1 < NUM_ATOMS && atom2 < NUM_ATOMS && atom1 != atom2 && r2 < CUTOFF_SQUARED);
#else #else
......
platforms/opencl/src/kernels/gbsaObc_nvidia.cl
View file @ b3dbaeaa
...@@ -108,7 +108,7 @@ __kernel void computeBornSum( ...@@ -108,7 +108,7 @@ __kernel void computeBornSum(
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
bornSum += l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) + bornSum += l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) +
(params2.y*params2.y*invR)*(l_ij2-u_ij2)); (params2.y*params2.y*invR)*(l_ij2-u_ij2));
if (params1.x < params2.x-r) if (params1.x < params2.y-r)
bornSum += 2.0f*(RECIP(params1.x)-l_ij); bornSum += 2.0f*(RECIP(params1.x)-l_ij);
} }
} }
...@@ -188,7 +188,7 @@ __kernel void computeBornSum( ...@@ -188,7 +188,7 @@ __kernel void computeBornSum(
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
float term = l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) + float term = l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) +
(params1.y*params1.y*invR)*(l_ij2-u_ij2)); (params1.y*params1.y*invR)*(l_ij2-u_ij2));
if (params2.x < params1.x-r) if (params2.x < params1.y-r)
term += 2.0f*(RECIP(params2.x)-l_ij); term += 2.0f*(RECIP(params2.x)-l_ij);
tempBuffer[get_local_id(0)] = term; tempBuffer[get_local_id(0)] = term;
} }
...@@ -247,7 +247,7 @@ __kernel void computeBornSum( ...@@ -247,7 +247,7 @@ __kernel void computeBornSum(
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
float term = l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) + float term = l_ij - u_ij + (0.50f*invR*ratio) + 0.25f*(r*(u_ij2-l_ij2) +
(params1.y*params1.y*invR)*(l_ij2-u_ij2)); (params1.y*params1.y*invR)*(l_ij2-u_ij2));
if (params2.x < params1.x-r) if (params2.x < params1.y-r)
term += 2.0f*(RECIP(params2.x)-l_ij); term += 2.0f*(RECIP(params2.x)-l_ij);
localData[tbx+tj].bornSum += term; localData[tbx+tj].bornSum += term;
} }
......
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