Commit 8868141e authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed compilation errors under AMD's OpenCL

parent 7482534e
const int TileSize = 32; #define TILE_SIZE 32
/** /**
* Find a bounding box for the atoms in each block. * Find a bounding box for the atoms in each block.
*/ */
__kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global float4* posq, __global float4* blockCenter, __global float4* blockBoundingBox) { __kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global float4* posq, __global float4* blockCenter, __global float4* blockBoundingBox) {
int index = get_global_id(0); int index = get_global_id(0);
int base = index*TileSize; int base = index*TILE_SIZE;
while (base < numAtoms) { while (base < numAtoms) {
float4 pos = posq[base]; float4 pos = posq[base];
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
...@@ -16,7 +16,7 @@ __kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global flo ...@@ -16,7 +16,7 @@ __kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global flo
#endif #endif
float4 minPos = pos; float4 minPos = pos;
float4 maxPos = pos; float4 maxPos = pos;
int last = min(base+TileSize, numAtoms); int last = min(base+TILE_SIZE, numAtoms);
for (int i = base+1; i < last; i++) { for (int i = base+1; i < last; i++) {
pos = posq[i]; pos = posq[i];
#ifdef USE_PERIODIC #ifdef USE_PERIODIC
...@@ -30,7 +30,7 @@ __kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global flo ...@@ -30,7 +30,7 @@ __kernel void findBlockBounds(int numAtoms, float4 periodicBoxSize, __global flo
blockBoundingBox[index] = 0.5f*(maxPos-minPos); blockBoundingBox[index] = 0.5f*(maxPos-minPos);
blockCenter[index] = 0.5f*(maxPos+minPos); blockCenter[index] = 0.5f*(maxPos+minPos);
index += get_global_size(0); index += get_global_size(0);
base = index*TileSize; base = index*TILE_SIZE;
} }
} }
...@@ -72,12 +72,12 @@ __kernel void findBlocksWithInteractions(int numTiles, float cutoffSquared, floa ...@@ -72,12 +72,12 @@ __kernel void findBlocksWithInteractions(int numTiles, float cutoffSquared, floa
*/ */
__kernel void findInteractionsWithinBlocks(float cutoffSquared, float4 periodicBoxSize, __global float4* posq, __global unsigned int* tiles, __global float4* blockCenter, __kernel void findInteractionsWithinBlocks(float cutoffSquared, float4 periodicBoxSize, __global float4* posq, __global unsigned int* tiles, __global float4* blockCenter,
__global float4* blockBoundingBox, __global unsigned int* interactionFlags, __global unsigned int* interactionCount, __local unsigned int* flags) { __global float4* blockBoundingBox, __global unsigned int* interactionFlags, __global unsigned int* interactionCount, __local unsigned int* flags) {
unsigned int totalWarps = get_global_size(0)/TileSize; unsigned int totalWarps = get_global_size(0)/TILE_SIZE;
unsigned int warp = get_global_id(0)/TileSize; unsigned int warp = get_global_id(0)/TILE_SIZE;
unsigned int numTiles = interactionCount[0]; unsigned int numTiles = interactionCount[0];
unsigned int pos = warp*numTiles/totalWarps; unsigned int pos = warp*numTiles/totalWarps;
unsigned int end = (warp+1)*numTiles/totalWarps; unsigned int end = (warp+1)*numTiles/totalWarps;
unsigned int index = get_local_id(0) & (TileSize - 1); unsigned int index = get_local_id(0) & (TILE_SIZE - 1);
unsigned int lasty = 0xFFFFFFFF; unsigned int lasty = 0xFFFFFFFF;
float4 apos; float4 apos;
...@@ -99,7 +99,7 @@ __kernel void findInteractionsWithinBlocks(float cutoffSquared, float4 periodicB ...@@ -99,7 +99,7 @@ __kernel void findInteractionsWithinBlocks(float cutoffSquared, float4 periodicB
float4 center = blockCenter[x]; float4 center = blockCenter[x];
float4 boxSize = blockBoundingBox[x]; float4 boxSize = blockBoundingBox[x];
if (y != lasty) if (y != lasty)
apos = posq[y*TileSize+index]; apos = posq[y*TILE_SIZE+index];
// Find the distance of the atom from the bounding box. // Find the distance of the atom from the bounding box.
......
const float dielectricOffset = 0.009f; #define DIELECTRIC_OFFSET 0.009f
const float probeRadius = 0.14f; #define PROBE_RADIUS 0.14f
const float surfaceAreaFactor = -6.0f*3.14159265358979323846f*0.0216f*1000.0f*0.4184f; #define SURFACE_AREA_FACTOR -170.351730667551f //-6.0f*3.14159265358979323846f*0.0216f*1000.0f*0.4184f;
/** /**
* Reduce the Born sums to compute the Born radii. * Reduce the Born sums to compute the Born radii.
...@@ -24,7 +24,7 @@ __kernel void reduceBornSum(int bufferSize, int numBuffers, float alpha, float b ...@@ -24,7 +24,7 @@ __kernel void reduceBornSum(int bufferSize, int numBuffers, float alpha, float b
float sum2 = sum*sum; float sum2 = sum*sum;
float sum3 = sum*sum2; float sum3 = sum*sum2;
float tanhSum = tanh(alpha*sum - beta*sum2 + gamma*sum3); float tanhSum = tanh(alpha*sum - beta*sum2 + gamma*sum3);
float nonOffsetRadius = offsetRadius + dielectricOffset; float nonOffsetRadius = offsetRadius + DIELECTRIC_OFFSET;
float radius = 1.0f/(1.0f/offsetRadius - tanhSum/nonOffsetRadius); float radius = 1.0f/(1.0f/offsetRadius - tanhSum/nonOffsetRadius);
float chain = offsetRadius*(alpha - 2.0f*beta*sum + 3.0f*gamma*sum2); float chain = offsetRadius*(alpha - 2.0f*beta*sum + 3.0f*gamma*sum2);
chain = (1.0f-tanhSum*tanhSum)*chain / nonOffsetRadius; chain = (1.0f-tanhSum*tanhSum)*chain / nonOffsetRadius;
...@@ -54,9 +54,9 @@ __kernel void reduceBornForce(int bufferSize, int numBuffers, __global float* bo ...@@ -54,9 +54,9 @@ __kernel void reduceBornForce(int bufferSize, int numBuffers, __global float* bo
float offsetRadius = params[index].x; float offsetRadius = params[index].x;
float bornRadius = bornRadii[index]; float bornRadius = bornRadii[index];
float r = offsetRadius+dielectricOffset+probeRadius; float r = offsetRadius+DIELECTRIC_OFFSET+PROBE_RADIUS;
float ratio6 = pow((offsetRadius+dielectricOffset)/bornRadius, 6.0f); float ratio6 = pow((offsetRadius+DIELECTRIC_OFFSET)/bornRadius, 6.0f);
float saTerm = surfaceAreaFactor*r*r*ratio6; float saTerm = SURFACE_AREA_FACTOR*r*r*ratio6;
force += saTerm/bornRadius; force += saTerm/bornRadius;
energy += saTerm; energy += saTerm;
force *= bornRadius*bornRadius*obcChain[index]; force *= bornRadius*bornRadius*obcChain[index];
......
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