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

Fixed errors in mixed precision mode

parent 820a6baa
...@@ -977,7 +977,7 @@ double OpenCLIntegrationUtilities::computeKineticEnergy(double timeShift) { ...@@ -977,7 +977,7 @@ double OpenCLIntegrationUtilities::computeKineticEnergy(double timeShift) {
timeShiftKernel.setArg<cl::Buffer>(0, context.getVelm().getDeviceBuffer()); timeShiftKernel.setArg<cl::Buffer>(0, context.getVelm().getDeviceBuffer());
timeShiftKernel.setArg<cl::Buffer>(1, context.getForce().getDeviceBuffer()); timeShiftKernel.setArg<cl::Buffer>(1, context.getForce().getDeviceBuffer());
if (context.getUseDoublePrecision() || context.getUseMixedPrecision()) if (context.getUseDoublePrecision())
timeShiftKernel.setArg<cl_double>(2, timeShift); timeShiftKernel.setArg<cl_double>(2, timeShift);
else else
timeShiftKernel.setArg<cl_float>(2, (cl_float) timeShift); timeShiftKernel.setArg<cl_float>(2, (cl_float) timeShift);
......
...@@ -5,7 +5,8 @@ __kernel void timeShiftVelocities(__global mixed4* restrict velm, __global const ...@@ -5,7 +5,8 @@ __kernel void timeShiftVelocities(__global mixed4* restrict velm, __global const
for (int index = get_global_id(0); index < NUM_ATOMS; index += get_global_size(0)) { for (int index = get_global_id(0); index < NUM_ATOMS; index += get_global_size(0)) {
mixed4 velocity = velm[index]; mixed4 velocity = velm[index];
if (velocity.w != 0.0) { if (velocity.w != 0.0) {
velocity.xyz += timeShift*force[index].xyz*velocity.w; mixed4 f = convert_mixed4(force[index]);
velocity.xyz += timeShift*f.xyz*velocity.w;
velm[index] = velocity; velm[index] = velocity;
} }
} }
......
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