         lastimpty = tmpx;
       }
       if (ival == lastival) lastimpty = -1e7f;            // Eliminate values which are not at value boundaries
-      if (lastimpty > maximpty) {
-        maximpty = lastimpty;
+      if (lastimpty < minimpty) {
+        minimpty = lastimpty;
         bestival = lastival;
       }
 
 #pragma unroll
-      for (h = 1; h < 32; h = h + h) {                    // Find the cumulative max impurity and corresponding ival
-        tmpx = __shfl_up(maximpty, h);
+      for (h = 1; h < 32; h = h + h) {                    // Find the cumulative min impurity and corresponding ival
+        tmpx = __shfl_up(minimpty, h);
         tmp = __shfl_up(bestival, h);
-        if (threadIdx.x >= h && tmpx > maximpty) {
-          maximpty = tmpx;
+        if (threadIdx.x >= h && tmpx < minimpty) {
+          minimpty = tmpx;
           bestival = tmp;
         }        
       }
-      maximpty = __shfl(maximpty, jtodo-1);               // Carefully copy the last active thread to all threads, needed outside this loop     
+      minimpty = __shfl(minimpty, jtodo-1);               // Carefully copy the last active thread to all threads, needed outside this loop     
       bestival = __shfl(bestival, jtodo-1);
       ctot = __shfl(ctot, jtodo-1);                
       clogc = __shfl(clogc, jtodo-1);

      ctot += cnt;                                        // Now update the total c and total ci log ci sums
      clogc += update;
      ctot = max(1, ctot);
      impty = (clogc) / ctot - log((float)ctot);          // And the impurity for this input

      tmp = __shfl_up(ival, 1);
      tmpx = __shfl_up(impty, 1);                         // Need the last impurity and ival in order


    lastival = -1;
    lastimpty = -1e7f;
    minimpty = -1e7f;
    ctot = 0;
    clogc = 0.0f;
    for (j = jc0; j < jc1; j += blockDim.x) {
