Skip to content

GitLab

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

Fix for OBC Born sum

parent f0cfae42
Show whitespace changes
Inline Side-by-side
Showing with 6 additions and 6 deletions
platforms/opencl/src/kernels/gbsaObc_cpu.cl
View file @ dbccf50f
...@@ -95,7 +95,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c ...@@ -95,7 +95,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
bornSum += l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + bornSum += l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2); (0.25f*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);
} }
} }
...@@ -146,7 +146,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c ...@@ -146,7 +146,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
bornSum += l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + bornSum += l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2); (0.25f*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);
} }
float rScaledRadiusI = r+params1.y; float rScaledRadiusI = r+params1.y;
...@@ -158,7 +158,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c ...@@ -158,7 +158,7 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
float term = l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + float term = l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params1.y*params1.y*invR)*(l_ij2-u_ij2); (0.25f*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[j].bornSum += term; localData[j].bornSum += term;
} }
......
platforms/opencl/src/kernels/gbsaObc_default.cl
View file @ dbccf50f
...@@ -90,7 +90,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa ...@@ -90,7 +90,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
bornSum += select(0.0f, l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + bornSum += select(0.0f, l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2), includeInteraction); (0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2), includeInteraction);
bornSum += select(0.0f, 2.0f*(RECIP(params1.x)-l_ij), includeInteraction && params1.x < params2.x-r); bornSum += select(0.0f, 2.0f*(RECIP(params1.x)-l_ij), includeInteraction && params1.x < params2.y-r);
} }
// Sum the forces and write results. // Sum the forces and write results.
...@@ -153,7 +153,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa ...@@ -153,7 +153,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa
unsigned int includeTerm = (includeInteraction && params1.x < rScaledRadiusJ); unsigned int includeTerm = (includeInteraction && params1.x < rScaledRadiusJ);
bornSum += select(0.0f, l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + bornSum += select(0.0f, l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2), includeTerm); (0.25f*params2.y*params2.y*invR)*(l_ij2-u_ij2), includeTerm);
bornSum += select(0.0f, 2.0f*(RECIP(params1.x)-l_ij), includeTerm && params1.x < params2.x-r); bornSum += select(0.0f, 2.0f*(RECIP(params1.x)-l_ij), includeTerm && params1.x < params2.y-r);
} }
float rScaledRadiusI = r+params1.y; float rScaledRadiusI = r+params1.y;
{ {
...@@ -164,7 +164,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa ...@@ -164,7 +164,7 @@ void computeBornSum(__global float* restrict global_bornSum, __global const floa
float ratio = LOG(u_ij * RECIP(l_ij)); float ratio = LOG(u_ij * RECIP(l_ij));
float term = l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) + float term = l_ij - u_ij + 0.25f*r*(u_ij2-l_ij2) + (0.50f*invR*ratio) +
(0.25f*params1.y*params1.y*invR)*(l_ij2-u_ij2); (0.25f*params1.y*params1.y*invR)*(l_ij2-u_ij2);
term += select(0.0f, 2.0f*(RECIP(params2.x)-l_ij), params2.x < params1.x-r); term += select(0.0f, 2.0f*(RECIP(params2.x)-l_ij), params2.x < params1.y-r);
localData[baseLocalAtom+tj+forceBufferOffset].bornSum += select(0.0f, term, includeInteraction && params2.x < rScaledRadiusI); localData[baseLocalAtom+tj+forceBufferOffset].bornSum += select(0.0f, term, includeInteraction && params2.x < rScaledRadiusI);
} }
barrier(CLK_LOCAL_MEM_FENCE); barrier(CLK_LOCAL_MEM_FENCE);
......
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