   long long key;
   int cnt, ival, icat, lastival;
   int ccnt, ctot, cnew, bestival, tmp;
-  float update, tmpx, clogc, impty, maximpty, lastimpty;
+  float update, tmpx, clogc, impty, minimpty, lastimpty;
 
   for (i = threadIdx.y + blockDim.y * blockIdx.x; i < ntrees*nnodes*nsamps; i += blockDim.y * gridDim.x) {
     // Process a group with fixed itree, inode, and ifeat

        lastimpty = tmpx;
      }
      if (ival == lastival) lastimpty = -1e7f;            // Eliminate values which are not at value boundaries
      if (lastimpty > maximpty) {
        maximpty = 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);
        tmp = __shfl_up(bestival, h);
        if (threadIdx.x >= h && tmpx > maximpty) {
          maximpty = tmpx;
          bestival = tmp;
        }        
      }
      maximpty = __shfl(maximpty, 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);

    }
    if (threadIdx.x == 0) {
      out[i] = bestival;                                  // Output the best split feature value
      outv[i] = minimpty + (clogc) / ctot - log((float)ctot);    // And the impurity gain
    }
  }
}
#else
__global__ void __minImpurity(long long *keys, int *counts, int *out, float *outv, int *jc, int *fieldlens, 
                              int ntrees, int nnodes, int ncats, int nsamps) {}
#endif

int minImpurity(long long *keys, int *counts, int *out, float *outv, int *jc, int *fieldlens, int ntrees, int nnodes, int ncats, int nsamps) {
  int ny = min(32, DBSIZE/ncats);
  dim3 tdim(32, ny, 1);
  int ng = min(64, 1L*ntrees*nnodes*nsamps);
  __minImpurity<<<ng,tdim>>>(keys, counts, out, outv, jc, fieldlens, ntrees, nnodes, ncats, nsamps);
  cudaDeviceSynchronize();
  cudaError_t err = cudaGetLastError();
  return err;
