Commit 600a7afa authored by Peter Eastman's avatar Peter Eastman
Browse files

Fixed bug in CPU version of GBSA

parent ad2d5c0a
...@@ -3,19 +3,16 @@ ...@@ -3,19 +3,16 @@
typedef struct { typedef struct {
float x, y, z; float x, y, z;
float q; float q;
float fx, fy, fz, fw;
float radius, scaledRadius; float radius, scaledRadius;
float bornSum; float bornSum;
float bornRadius; } AtomData1;
float bornForce;
} AtomData;
/** /**
* Compute the Born sum. * Compute the Born sum.
*/ */
__kernel void computeBornSum(__global float* restrict global_bornSum, __global const float4* restrict posq, __global const float2* restrict global_params, __kernel void computeBornSum(__global float* restrict global_bornSum, __global const float4* restrict posq, __global const float2* restrict global_params,
__local AtomData* restrict localData, __local AtomData1* restrict localData,
#ifdef USE_CUTOFF #ifdef USE_CUTOFF
__global const ushort2* restrict tiles, __global const unsigned int* restrict interactionCount, float4 periodicBoxSize, float4 invPeriodicBoxSize, unsigned int maxTiles, __global const unsigned int* restrict interactionFlags) { __global const ushort2* restrict tiles, __global const unsigned int* restrict interactionCount, float4 periodicBoxSize, float4 invPeriodicBoxSize, unsigned int maxTiles, __global const unsigned int* restrict interactionFlags) {
#else #else
...@@ -186,13 +183,20 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c ...@@ -186,13 +183,20 @@ __kernel void computeBornSum(__global float* restrict global_bornSum, __global c
} }
} }
typedef struct {
float x, y, z;
float q;
float fx, fy, fz, fw;
float bornRadius;
} AtomData2;
/** /**
* First part of computing the GBSA interaction. * First part of computing the GBSA interaction.
*/ */
__kernel void computeGBSAForce1(__global float4* restrict forceBuffers, __global float* restrict energyBuffer, __kernel void computeGBSAForce1(__global float4* restrict forceBuffers, __global float* restrict global_bornForce,
__global const float4* restrict posq, __global const float* restrict global_bornRadii, __global float* restrict global_bornForce, __global float* restrict energyBuffer, __global const float4* restrict posq, __global const float* restrict global_bornRadii,
__local AtomData* restrict localData, __local AtomData2* restrict localData,
#ifdef USE_CUTOFF #ifdef USE_CUTOFF
__global const ushort2* restrict tiles, __global const unsigned int* restrict interactionCount, float4 periodicBoxSize, float4 invPeriodicBoxSize, unsigned int maxTiles, __global const unsigned int* restrict interactionFlags) { __global const ushort2* restrict tiles, __global const unsigned int* restrict interactionCount, float4 periodicBoxSize, float4 invPeriodicBoxSize, unsigned int maxTiles, __global const unsigned int* restrict interactionFlags) {
#else #else
......
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