Skip to content

GitLab

Commit a8ceded2 authored by Peter Eastman's avatar Peter Eastman
Browse files

Reverted the changes from revision 2214, which triggered a bug in the ATI compiler

parent 1763a764
Show whitespace changes
Inline Side-by-side
Showing with 51 additions and 53 deletions
platforms/opencl/src/OpenCLNonbondedUtilities.cpp
View file @ a8ceded2
......@@ -325,23 +325,23 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
stringstream loadLocal1;
for (int i = 0; i < (int) params.size(); i++) {
if (params[i].getNumComponents() == 1) {
loadLocal1<<"atom1Data->"<<params[i].getName()<<" = "<<params[i].getName()<<"1;\n";
loadLocal1<<"localData[get_local_id(0)]."<<params[i].getName()<<" = "<<params[i].getName()<<"1;\n";
}
else {
for (int j = 0; j < params[i].getNumComponents(); ++j)
loadLocal1<<"atom1Data->"<<params[i].getName()<<"_"<<suffixes[j]<<" = "<<params[i].getName()<<"1."<<suffixes[j]<<";\n";
loadLocal1<<"localData[get_local_id(0)]."<<params[i].getName()<<"_"<<suffixes[j]<<" = "<<params[i].getName()<<"1."<<suffixes[j]<<";\n";
}
}
replacements["LOAD_LOCAL_PARAMETERS_FROM_1"] = loadLocal1.str();
stringstream loadLocal2;
for (int i = 0; i < (int) params.size(); i++) {
if (params[i].getNumComponents() == 1) {
loadLocal2<<"atom1Data->"<<params[i].getName()<<" = global_"<<params[i].getName()<<"[j];\n";
loadLocal2<<"localData[get_local_id(0)]."<<params[i].getName()<<" = global_"<<params[i].getName()<<"[j];\n";
}
else {
loadLocal2<<params[i].getType()<<" temp_"<<params[i].getName()<<" = global_"<<params[i].getName()<<"[j];\n";
for (int j = 0; j < params[i].getNumComponents(); ++j)
loadLocal2<<"atom1Data->"<<params[i].getName()<<"_"<<suffixes[j]<<" = temp_"<<params[i].getName()<<"."<<suffixes[j]<<";\n";
loadLocal2<<"localData[get_local_id(0)]."<<params[i].getName()<<"_"<<suffixes[j]<<" = temp_"<<params[i].getName()<<"."<<suffixes[j]<<";\n";
}
}
replacements["LOAD_LOCAL_PARAMETERS_FROM_GLOBAL"] = loadLocal2.str();
......@@ -358,14 +358,14 @@ cl::Kernel OpenCLNonbondedUtilities::createInteractionKernel(const string& sourc
stringstream load2j;
for (int i = 0; i < (int) params.size(); i++) {
if (params[i].getNumComponents() == 1) {
load2j<<params[i].getType()<<" "<<params[i].getName()<<"2 = atom2Data->"<<params[i].getName()<<";\n";
load2j<<params[i].getType()<<" "<<params[i].getName()<<"2 = localData[atom2]."<<params[i].getName()<<";\n";
}
else {
load2j<<params[i].getType()<<" "<<params[i].getName()<<"2 = ("<<params[i].getType()<<") (";
for (int j = 0; j < params[i].getNumComponents(); ++j) {
if (j > 0)
load2j<<", ";
load2j<<"atom2Data->"<<params[i].getName()<<"_"<<suffixes[j];
load2j<<"localData[atom2]."<<params[i].getName()<<"_"<<suffixes[j];
}
load2j<<");\n";
}
......
platforms/opencl/src/kernels/nonbonded_default.cl
View file @ a8ceded2
......@@ -40,15 +40,14 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
unsigned int atom1 = x + tgx;
float4 force = 0.0f;
float4 posq1 = posq[atom1];
__local AtomData* atom1Data = &localData[get_local_id(0)];
LOAD_ATOM1_PARAMETERS
if (x == y) {
// This tile is on the diagonal.
atom1Data->x = posq1.x;
atom1Data->y = posq1.y;
atom1Data->z = posq1.z;
atom1Data->q = posq1.w;
localData[get_local_id(0)].x = posq1.x;
localData[get_local_id(0)].y = posq1.y;
localData[get_local_id(0)].z = posq1.z;
localData[get_local_id(0)].q = posq1.w;
LOAD_LOCAL_PARAMETERS_FROM_1
barrier(CLK_LOCAL_MEM_FENCE);
unsigned int xi = x/TILE_SIZE;
......@@ -60,8 +59,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
#ifdef USE_EXCLUSIONS
bool isExcluded = !(excl & 0x1);
#endif
__local AtomData* atom2Data = &localData[baseLocalAtom+j];
float4 posq2 = (float4) (atom2Data->x, atom2Data->y, atom2Data->z, atom2Data->q);
int atom2 = baseLocalAtom+j;
float4 posq2 = (float4) (localData[atom2].x, localData[atom2].y, localData[atom2].z, localData[atom2].q);
float4 delta = (float4) (posq2.xyz - posq1.xyz, 0.0f);
#ifdef USE_PERIODIC
delta.x -= floor(delta.x*INV_PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
......@@ -72,7 +71,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
float invR = RSQRT(r2);
float r = RECIP(invR);
LOAD_ATOM2_PARAMETERS
int atom2 = y+baseLocalAtom+j;
atom2 = y+baseLocalAtom+j;
float dEdR = 0.0f;
float tempEnergy = 0.0f;
COMPUTE_INTERACTION
......@@ -102,15 +101,15 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
if (lasty != y && get_local_id(0) < TILE_SIZE) {
unsigned int j = y + tgx;
float4 tempPosq = posq[j];
atom1Data->x = tempPosq.x;
atom1Data->y = tempPosq.y;
atom1Data->z = tempPosq.z;
atom1Data->q = tempPosq.w;
localData[get_local_id(0)].x = tempPosq.x;
localData[get_local_id(0)].y = tempPosq.y;
localData[get_local_id(0)].z = tempPosq.z;
localData[get_local_id(0)].q = tempPosq.w;
LOAD_LOCAL_PARAMETERS_FROM_GLOBAL
}
atom1Data->fx = 0.0f;
atom1Data->fy = 0.0f;
atom1Data->fz = 0.0f;
localData[get_local_id(0)].fx = 0.0f;
localData[get_local_id(0)].fy = 0.0f;
localData[get_local_id(0)].fz = 0.0f;
barrier(CLK_LOCAL_MEM_FENCE);
// Compute the full set of interactions in this tile.
......@@ -128,8 +127,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
#ifdef USE_EXCLUSIONS
bool isExcluded = !(excl & 0x1);
#endif
__local AtomData* atom2Data = &localData[baseLocalAtom+j];
float4 posq2 = (float4) (atom2Data->x, atom2Data->y, atom2Data->z, atom2Data->q);
int atom2 = baseLocalAtom+tj;
float4 posq2 = (float4) (localData[atom2].x, localData[atom2].y, localData[atom2].z, localData[atom2].q);
float4 delta = (float4) (posq2.xyz - posq1.xyz, 0.0f);
#ifdef USE_PERIODIC
delta.x -= floor(delta.x*INV_PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
......@@ -140,7 +139,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
float invR = RSQRT(r2);
float r = RECIP(invR);
LOAD_ATOM2_PARAMETERS
int atom2 = y+baseLocalAtom+tj;
atom2 = y+baseLocalAtom+tj;
float dEdR = 0.0f;
float tempEnergy = 0.0f;
COMPUTE_INTERACTION
......@@ -169,7 +168,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
unsigned int offset2 = y + tgx + get_group_id(0)*PADDED_NUM_ATOMS;
#endif
forceBuffers[offset1].xyz = forceBuffers[offset1].xyz+force.xyz+tempBuffer[get_local_id(0)+TILE_SIZE].xyz;
float4 sum = (float4) (atom1Data->fx+localData[get_local_id(0)+TILE_SIZE].fx, atom1Data->fy+localData[get_local_id(0)+TILE_SIZE].fy, atom1Data->fz+localData[get_local_id(0)+TILE_SIZE].fz, 0.0f);
float4 sum = (float4) (localData[get_local_id(0)].fx+localData[get_local_id(0)+TILE_SIZE].fx, localData[get_local_id(0)].fy+localData[get_local_id(0)+TILE_SIZE].fy, localData[get_local_id(0)].fz+localData[get_local_id(0)+TILE_SIZE].fz, 0.0f);
forceBuffers[offset2].xyz = forceBuffers[offset2].xyz+sum.xyz;
}
lasty = y;
......
platforms/opencl/src/kernels/nonbonded_nvidia.cl
View file @ a8ceded2
......@@ -41,15 +41,14 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
unsigned int atom1 = x + tgx;
float4 force = 0.0f;
float4 posq1 = posq[atom1];
__local AtomData* atom1Data = &localData[get_local_id(0)];
LOAD_ATOM1_PARAMETERS
if (x == y) {
// This tile is on the diagonal.
atom1Data->x = posq1.x;
atom1Data->y = posq1.y;
atom1Data->z = posq1.z;
atom1Data->q = posq1.w;
localData[get_local_id(0)].x = posq1.x;
localData[get_local_id(0)].y = posq1.y;
localData[get_local_id(0)].z = posq1.z;
localData[get_local_id(0)].q = posq1.w;
LOAD_LOCAL_PARAMETERS_FROM_1
unsigned int xi = x/TILE_SIZE;
unsigned int tile = xi+xi*PADDED_NUM_ATOMS/TILE_SIZE-xi*(xi+1)/2;
......@@ -60,8 +59,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
#ifdef USE_EXCLUSIONS
bool isExcluded = !(excl & 0x1);
#endif
__local AtomData* atom2Data = &localData[tbx+j];
float4 posq2 = (float4) (atom2Data->x, atom2Data->y, atom2Data->z, atom2Data->q);
int atom2 = tbx+j;
float4 posq2 = (float4) (localData[atom2].x, localData[atom2].y, localData[atom2].z, localData[atom2].q);
float4 delta = (float4) (posq2.xyz - posq1.xyz, 0.0f);
#ifdef USE_PERIODIC
delta.x -= floor(delta.x*INV_PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
......@@ -72,7 +71,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
float r = sqrt(r2);
float invR = 1.0f/r;
LOAD_ATOM2_PARAMETERS
int atom2 = y+j;
atom2 = y+j;
float dEdR = 0.0f;
float tempEnergy = 0.0f;
COMPUTE_INTERACTION
......@@ -96,15 +95,15 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
if (lasty != y) {
unsigned int j = y + tgx;
float4 tempPosq = posq[j];
atom1Data->x = tempPosq.x;
atom1Data->y = tempPosq.y;
atom1Data->z = tempPosq.z;
atom1Data->q = tempPosq.w;
localData[get_local_id(0)].x = tempPosq.x;
localData[get_local_id(0)].y = tempPosq.y;
localData[get_local_id(0)].z = tempPosq.z;
localData[get_local_id(0)].q = tempPosq.w;
LOAD_LOCAL_PARAMETERS_FROM_GLOBAL
}
atom1Data->fx = 0.0f;
atom1Data->fy = 0.0f;
atom1Data->fz = 0.0f;
localData[get_local_id(0)].fx = 0.0f;
localData[get_local_id(0)].fy = 0.0f;
localData[get_local_id(0)].fz = 0.0f;
#ifdef USE_CUTOFF
unsigned int flags = interactionFlags[pos];
if (!hasExclusions && flags != 0xFFFFFFFF) {
......@@ -117,8 +116,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
for (unsigned int j = 0; j < TILE_SIZE; j++) {
if ((flags&(1<<j)) != 0) {
bool isExcluded = false;
__local AtomData* atom2Data = &localData[tbx+j];
float4 posq2 = (float4) (atom2Data->x, atom2Data->y, atom2Data->z, atom2Data->q);
int atom2 = tbx+j;
float4 posq2 = (float4) (localData[atom2].x, localData[atom2].y, localData[atom2].z, localData[atom2].q);
float4 delta = (float4) (posq2.xyz - posq1.xyz, 0.0f);
#ifdef USE_PERIODIC
delta.x -= floor(delta.x*INV_PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
......@@ -129,7 +128,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
float invR = RSQRT(r2);
float r = RECIP(invR);
LOAD_ATOM2_PARAMETERS
int atom2 = y+j;
atom2 = y+j;
float dEdR = 0.0f;
float tempEnergy = 0.0f;
COMPUTE_INTERACTION
......@@ -149,9 +148,9 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
if (tgx % 16 == 0)
tempBuffer[get_local_id(0)].xyz += tempBuffer[get_local_id(0)+8].xyz;
if (tgx == 0) {
atom2Data->fx += tempBuffer[get_local_id(0)].x + tempBuffer[get_local_id(0)+16].x;
atom2Data->fy += tempBuffer[get_local_id(0)].y + tempBuffer[get_local_id(0)+16].y;
atom2Data->fz += tempBuffer[get_local_id(0)].z + tempBuffer[get_local_id(0)+16].z;
localData[tbx+j].fx += tempBuffer[get_local_id(0)].x + tempBuffer[get_local_id(0)+16].x;
localData[tbx+j].fy += tempBuffer[get_local_id(0)].y + tempBuffer[get_local_id(0)+16].y;
localData[tbx+j].fz += tempBuffer[get_local_id(0)].z + tempBuffer[get_local_id(0)+16].z;
}
}
}
......@@ -174,8 +173,8 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
#ifdef USE_EXCLUSIONS
bool isExcluded = !(excl & 0x1);
#endif
__local AtomData* atom2Data = &localData[tbx+tj];
float4 posq2 = (float4) (atom2Data->x, atom2Data->y, atom2Data->z, atom2Data->q);
int atom2 = tbx+tj;
float4 posq2 = (float4) (localData[atom2].x, localData[atom2].y, localData[atom2].z, localData[atom2].q);
float4 delta = (float4) (posq2.xyz - posq1.xyz, 0.0f);
#ifdef USE_PERIODIC
delta.x -= floor(delta.x*INV_PERIODIC_BOX_SIZE_X+0.5f)*PERIODIC_BOX_SIZE_X;
......@@ -186,16 +185,16 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
float invR = RSQRT(r2);
float r = RECIP(invR);
LOAD_ATOM2_PARAMETERS
int atom2 = y+tj;
atom2 = y+tj;
float dEdR = 0.0f;
float tempEnergy = 0.0f;
COMPUTE_INTERACTION
energy += tempEnergy;
delta.xyz *= dEdR;
force.xyz -= delta.xyz;
atom2Data->fx += delta.x;
atom2Data->fy += delta.y;
atom2Data->fz += delta.z;
localData[tbx+tj].fx += delta.x;
localData[tbx+tj].fy += delta.y;
localData[tbx+tj].fz += delta.z;
excl >>= 1;
tj = (tj + 1) & (TILE_SIZE - 1);
}
......@@ -210,7 +209,7 @@ void computeNonbonded(__global float4* forceBuffers, __global float* energyBuffe
unsigned int offset2 = y + tgx + warp*PADDED_NUM_ATOMS;
#endif
forceBuffers[offset1].xyz += force.xyz;
forceBuffers[offset2] += (float4) (atom1Data->fx, atom1Data->fy, atom1Data->fz, 0.0f);
forceBuffers[offset2] += (float4) (localData[get_local_id(0)].fx, localData[get_local_id(0)].fy, localData[get_local_id(0)].fz, 0.0f);
lasty = y;
}
pos++;
......
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