Workaround for compiler bug on Lion

platforms/opencl/src/OpenCLKernels.cpp

@@ -1530,6 +1530,9 @@ double OpenCLCalcGBSAOBCForceKernel::execute(ContextImpl& context, bool includeF
         defines["PADDED_NUM_ATOMS"] = intToString(cl.getPaddedNumAtoms());
         defines["NUM_BLOCKS"] = OpenCLExpressionUtilities::intToString(cl.getNumAtomBlocks());
         defines["FORCE_WORK_GROUP_SIZE"] = OpenCLExpressionUtilities::intToString(nb.getForceThreadBlockSize());
+        string platformVendor = cl::Platform(cl.getDevice().getInfo<CL_DEVICE_PLATFORM>()).getInfo<CL_PLATFORM_VENDOR>();
+        if (platformVendor == "Apple")
+            defines["USE_APPLE_WORKAROUND"] = "1";
         string file;
         if (deviceIsCpu)
             file = OpenCLKernelSources::gbsaObc_cpu;

platforms/opencl/src/kernels/gbsaObc_nvidia.cl

@@ -457,6 +457,9 @@ __kernel void computeGBSAForce1(
 #ifdef USE_CUTOFF
                 unsigned int flags = (numTiles <= maxTiles ? interactionFlags[pos] : 0xFFFFFFFF);
                 bool computeSubset = false;
+#ifdef USE_APPLE_WORKAROUND
+                computeSubset = (flags == 0); // Workaround for a compiler bug in Apple's OpenCL on Lion
+#else
                 if (flags != 0xFFFFFFFF) {
                     if (tgx < 2)
                         exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
@@ -467,6 +470,7 @@ __kernel void computeGBSAForce1(
                             exclusionIndex[localGroupIndex] = i*TILE_SIZE;
                     computeSubset = (exclusionIndex[localGroupIndex] == -1);
                 }
+#endif
                 if (computeSubset) {
                     if (flags == 0) {
                         // No interactions in this tile.