 
 #if __CUDA_ARCH__ >= 300
 
-__global__ void __maxImpurity(long long *keys, int *counts, int *out, float *outv, int *jc, int *fieldlens, 
+__global__ void __minImpurity(long long *keys, int *counts, int *out, float *outv, int *jc, int *fieldlens, 
                               int ntrees, int nnodes, int ncats, int nsamps) {
   __shared__ int catcnt[DBSIZE];
   __shared__ int fl[32];

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

int maxImpurity(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);
  __maxImpurity<<<ng,tdim>>>(keys, counts, out, outv, jc, fieldlens, ntrees, nnodes, ncats, nsamps);
  cudaDeviceSynchronize();
  cudaError_t err = cudaGetLastError();
  return err;

        lastimpty = tmpx;
      }
      if (ival == lastival) lastimpty = -1e7f;            // Eliminate values which are not at value boundaries
      if (lastimpty < minimpty) {
        minimpty = lastimpty;
        bestival = lastival;
      }

#pragma unroll
      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 < minimpty) {
          minimpty = tmpx;
          bestival = tmp;
        }        
      }
      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);
