Fixed errors under Intel OpenCL

platforms/opencl/src/OpenCLKernels.cpp

@@ -3517,7 +3517,7 @@ void OpenCLCalcCustomHbondForceKernel::initialize(const System& system, const Cu
     defines["PADDED_NUM_ATOMS"] = intToString(cl.getPaddedNumAtoms());
     defines["NUM_DONORS"] = intToString(numDonors);
     defines["NUM_ACCEPTORS"] = intToString(numAcceptors);
-    defines["M_PI"] = doubleToString(M_PI);
+    defines["PI"] = doubleToString(M_PI);
     if (force.getNonbondedMethod() != CustomHbondForce::NoCutoff) {
         defines["USE_CUTOFF"] = "1";
         defines["CUTOFF_SQUARED"] = doubleToString(force.getCutoffDistance()*force.getCutoffDistance());

platforms/opencl/src/OpenCLNonbondedUtilities.cpp

@@ -375,10 +375,12 @@ void OpenCLNonbondedUtilities::updateNeighborListSize() {
         interactionFlags = new OpenCLArray<cl_uint>(context, deviceIsCpu ? 2*newSize : newSize, "interactionFlags");
         forceKernel.setArg<cl::Buffer>(13, interactionFlags->getDeviceBuffer());
         findInteractingBlocksKernel.setArg<cl::Buffer>(7, interactionFlags->getDeviceBuffer());
+		if (!deviceIsCpu) {
             findInteractionsWithinBlocksKernel.setArg<cl::Buffer>(4, interactingTiles->getDeviceBuffer());
             findInteractionsWithinBlocksKernel.setArg<cl::Buffer>(7, interactionFlags->getDeviceBuffer());
             findInteractionsWithinBlocksKernel.setArg<cl_uint>(10, newSize);
 		}
+    }
 }
 
 void OpenCLNonbondedUtilities::setTileRange(int startTileIndex, int numTiles) {

platforms/opencl/src/kernels/customHbondForce.cl

@@ -36,7 +36,7 @@ float computeAngle(float4 vec1, float4 vec2) {
         float scale = vec1.w*vec2.w;
         angle = asin(SQRT(dot(crossProduct, crossProduct)/scale));
         if (cosine < 0.0f)
-            angle = M_PI-angle;
+            angle = PI-angle;
     }
     else
        angle = acos(cosine);

platforms/opencl/src/kernels/sort.cl

@@ -30,6 +30,7 @@ __kernel void computeRange(__global const DATA_TYPE* restrict data, uint length,
         barrier(CLK_LOCAL_MEM_FENCE);
     }
     minimum = buffer[0];
+    barrier(CLK_LOCAL_MEM_FENCE);
     buffer[get_local_id(0)] = maximum;
     barrier(CLK_LOCAL_MEM_FENCE);
     for (uint step = 1; step < get_local_size(0); step *= 2) {