Fixed lots of warnings and errors on Lion

platforms/opencl/src/OpenCLContext.cpp

@@ -57,6 +57,9 @@ const int OpenCLContext::ThreadBlockSize = 64;
 const int OpenCLContext::TileSize = 32;
 
 static void CL_CALLBACK errorCallback(const char* errinfo, const void* private_info, size_t cb, void* user_data) {
+    string skip = "OpenCL Build Warning : Compiler build log:";
+    if (strncmp(errinfo, skip.c_str(), skip.length()) == 0)
+        return; // OS X Lion insists on calling this for every build warning, even though they aren't errors.
     std::cerr << "OpenCL internal error: " << errinfo << std::endl;
 }
 

platforms/opencl/src/kernels/customGBEnergyN2_nvidia.cl

@@ -79,7 +79,7 @@ __kernel void computeN2Energy(
                 exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
             if (tgx == 0)
                 exclusionIndex[localGroupIndex] = -1;
-            for (int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
+            for (unsigned int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
                 if (exclusionIndices[i] == y)
                     exclusionIndex[localGroupIndex] = i*TILE_SIZE;
             bool hasExclusions = (exclusionIndex[localGroupIndex] > -1);

platforms/opencl/src/kernels/customGBValueN2_nvidia.cl

@@ -72,7 +72,7 @@ __kernel void computeN2Value(__global const float4* restrict posq, __local float
                 exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
             if (tgx == 0)
                 exclusionIndex[localGroupIndex] = -1;
-            for (int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
+            for (unsigned int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
                 if (exclusionIndices[i] == y)
                     exclusionIndex[localGroupIndex] = i*TILE_SIZE;
             bool hasExclusions = (exclusionIndex[localGroupIndex] > -1);

platforms/opencl/src/kernels/gbsaObc_nvidia.cl

@@ -137,7 +137,7 @@ __kernel void computeBornSum(
                         exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
                     if (tgx == 0)
                         exclusionIndex[localGroupIndex] = -1;
-                    for (int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
+                    for (unsigned int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
                         if (exclusionIndices[i] == y)
                             exclusionIndex[localGroupIndex] = i*TILE_SIZE;
                     computeSubset = (exclusionIndex[localGroupIndex] == -1);
@@ -462,7 +462,7 @@ __kernel void computeGBSAForce1(
                         exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
                     if (tgx == 0)
                         exclusionIndex[localGroupIndex] = -1;
-                    for (int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
+                    for (unsigned int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
                         if (exclusionIndices[i] == y)
                             exclusionIndex[localGroupIndex] = i*TILE_SIZE;
                     computeSubset = (exclusionIndex[localGroupIndex] == -1);

platforms/opencl/src/kernels/langevin.cl

@@ -75,7 +75,7 @@ __kernel void selectLangevinStepSize(float maxStepSize, float errorTol, float ta
 
     // Sum the errors from all threads.
 
-    for (int offset = 1; offset < get_local_size(0); offset *= 2) {
+    for (unsigned int offset = 1; offset < get_local_size(0); offset *= 2) {
         if (get_local_id(0)+offset < get_local_size(0) && (get_local_id(0)&(2*offset-1)) == 0)
             error[get_local_id(0)] += error[get_local_id(0)+offset];
         barrier(CLK_LOCAL_MEM_FENCE);

platforms/opencl/src/kernels/nonbonded_nvidia.cl

@@ -81,7 +81,7 @@ __kernel void computeNonbonded(
                 exclusionRange[2*localGroupIndex+tgx] = exclusionRowIndices[x+tgx];
             if (tgx == 0)
                 exclusionIndex[localGroupIndex] = -1;
-            for (int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
+            for (unsigned int i = exclusionRange[2*localGroupIndex]+tgx; i < exclusionRange[2*localGroupIndex+1]; i += TILE_SIZE)
                 if (exclusionIndices[i] == y)
                     exclusionIndex[localGroupIndex] = i*TILE_SIZE;
             bool hasExclusions = (exclusionIndex[localGroupIndex] > -1);

platforms/opencl/src/kernels/removeCM.cl

@@ -53,10 +53,10 @@ __kernel void calcCenterOfMassMomentum(int numAtoms, __global const float4* rest
  * Remove center of mass motion.
  */
 
-__kernel void removeCenterOfMassMomentum(int numAtoms, __global float4* restrict velm, __global const float4* restrict cmMomentum, __local volatile float4* restrict temp) {
+__kernel void removeCenterOfMassMomentum(unsigned int numAtoms, __global float4* restrict velm, __global const float4* restrict cmMomentum, __local volatile float4* restrict temp) {
     // First sum all of the momenta that were calculated by individual groups.
 
-    int index = get_local_id(0);
+    unsigned int index = get_local_id(0);
     float4 cm = 0.0f;
     while (index < get_num_groups(0)) {
         cm += cmMomentum[index];

platforms/opencl/src/kernels/verlet.cl

@@ -61,7 +61,7 @@ __kernel void selectVerletStepSize(int numAtoms, float maxStepSize, float errorT
     // Calculate the error.
 
     float err = 0.0f;
-    unsigned int index = get_local_id(0);
+    int index = get_local_id(0);
     while (index < numAtoms) {
         float4 f = force[index];
         float invMass = velm[index].w;
@@ -73,7 +73,7 @@ __kernel void selectVerletStepSize(int numAtoms, float maxStepSize, float errorT
 
     // Sum the errors from all threads.
 
-    for (int offset = 1; offset < get_local_size(0); offset *= 2) {
+    for (unsigned int offset = 1; offset < get_local_size(0); offset *= 2) {
         if (get_local_id(0)+offset < get_local_size(0) && (get_local_id(0)&(2*offset-1)) == 0)
             error[get_local_id(0)] += error[get_local_id(0)+offset];
         barrier(CLK_LOCAL_MEM_FENCE);

platforms/opencl/tests/TestOpenCLCustomAngleForce.cpp

@@ -123,15 +123,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, i%2, 0);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLCustomBondForce.cpp

@@ -130,15 +130,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, 0, 0);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLCustomExternalForce.cpp

@@ -130,15 +130,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(5*genrand_real2(sfmt), 5*genrand_real2(sfmt), 5*genrand_real2(sfmt));
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLCustomNonbondedForce.cpp

@@ -380,15 +380,14 @@ void testParallelComputation() {
                 force->addExclusion(i, j);
         }
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLCustomTorsionForce.cpp

@@ -162,15 +162,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, i%2, i%3);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLHarmonicAngleForce.cpp

@@ -92,15 +92,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, i%2, 0);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLHarmonicBondForce.cpp

@@ -87,15 +87,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, 0, 0);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLNonbondedForce.cpp

@@ -698,15 +698,14 @@ void testParallelComputation() {
                 force->addException(i, j, 0, 1, 0);
         }
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLPeriodicTorsionForce.cpp

@@ -89,15 +89,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, i%2, i%3);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);

platforms/opencl/tests/TestOpenCLRBTorsionForce.cpp

@@ -99,15 +99,14 @@ void testParallelComputation() {
     for (int i = 0; i < numParticles; i++)
         positions[i] = Vec3(i, i%2, i%3);
     VerletIntegrator integrator1(0.01);
-    map<string, string> props1;
-    props1[OpenCLPlatform::OpenCLDeviceIndex()] = "0";
-    Context context1(system, integrator1, platform, props1);
+    Context context1(system, integrator1, platform);
     context1.setPositions(positions);
     State state1 = context1.getState(State::Forces | State::Energy);
     VerletIntegrator integrator2(0.01);
-    map<string, string> props2;
-    props2[OpenCLPlatform::OpenCLDeviceIndex()] = "0,0";
-    Context context2(system, integrator2, platform, props2);
+    string deviceIndex = platform.getPropertyValue(context1, OpenCLPlatform::OpenCLDeviceIndex());
+    map<string, string> props;
+    props[OpenCLPlatform::OpenCLDeviceIndex()] = deviceIndex+","+deviceIndex;
+    Context context2(system, integrator2, platform, props);
     context2.setPositions(positions);
     State state2 = context2.getState(State::Forces | State::Energy);
     ASSERT_EQUAL_TOL(state1.getPotentialEnergy(), state2.getPotentialEnergy(), 1e-5);