Merge pull request #2553 from peastman/index

Fixed incorrect indexing in kernel

platforms/common/src/kernels/customNonbondedGroups.cc

@@ -23,7 +23,7 @@ DEVICE int reduceMax(int val, LOCAL_ARG int* temp) {
     temp[LOCAL_ID] = val;
     SYNC_WARPS;
     for (int offset = 16; offset > 0; offset /= 2) {
-        if (offset < indexInWarp)
+        if (indexInWarp < offset)
             temp[LOCAL_ID] = max(temp[LOCAL_ID], temp[LOCAL_ID+offset]);
         SYNC_WARPS;
     }

platforms/cuda/src/CudaContext.cpp

@@ -258,7 +258,7 @@ CudaContext::CudaContext(const System& system, int deviceIndex, bool useBlocking
     int multiprocessors;
     CHECK_RESULT(cuDeviceGetAttribute(&multiprocessors, CU_DEVICE_ATTRIBUTE_MULTIPROCESSOR_COUNT, device));
     numThreadBlocks = numThreadBlocksPerComputeUnit*multiprocessors;
-    if (computeCapability >= 7.0) {
+    if (cudaDriverVersion >= 9000) {
         compilationDefines["SYNC_WARPS"] = "__syncwarp();";
         compilationDefines["SHFL(var, srcLane)"] = "__shfl_sync(0xffffffff, var, srcLane);";
         compilationDefines["BALLOT(var)"] = "__ballot_sync(0xffffffff, var);";