+#define MINUS_LOG_THRESHOLD -18.42
+#define SOFTMAXTREE_THREADS 128
+
+__global__ void cunnx_SoftMaxTree_updateOutput_kernel(
+  float *output, float *input, float* weight, float* bias, 
+  int* meta, int nNode, int nInput)
+{
+  __shared__ float input_buffer[nInput];
+  int k = blockIdx.x;
+  float *input_k = input + k*dim;
+  float *output_k = output + k*dim;
+  
+  // copy input into buffer
+  for (int i=0; i<nInput ;i++) 
+  {
+    input_buffer[i] = input_k[i];
+  }
+
+  
+  while (1) 
+  {
+    float z = input_k[i];
+    if(buffer[threadIdx.x] < z)
+      buffer[threadIdx.x] = z;
+  }
+
+}
 
 static int cunnx_SoftMaxTree_updateOutput(lua_State *L)
 { 

  THCudaTensor *output = luaT_getfieldcheckudata(L, 1, "output", "torch.CudaTensor");
  
  THIntTensor *node;
  THCudaTensor *nodeWeight, *nodeBias, *nodeOutput, *nodeInput, *nodeInter;
  real *input_data, *output_data;

  long i, d;
  long n = 0;
  
  luaL_argcheck(L, input->nDimension == 2, 2, "2D(batch mode) tensor expected");
  luaL_argcheck(L, input->size[1] == inputSize, 2, "invalid input size");

  node = THIntTensor_new();
  nodeWeight = THTensor_(new)();
  nodeBias = THCudaTensor_(new)();
  nodeOutput = THCudaTensor_(new)();
  nodeInput = THCudaTensor_(new)();
  nodeInter = THCudaTensor_(new)();
  
  
  cudaError errcode = cudaGetLastError();
  if(errcode != cudaSuccess)

  THIntTensor *childParent = (THIntTensor*)luaT_getfieldcheckudata(L, 1, "childParent", "torch.IntTensor");
  THIntTensor *parentChildren = (THIntTensor*)luaT_getfieldcheckudata(L, 1, "parentChildren", "torch.IntTensor");
  
  THCudaTensor *metadata = luaT_getfieldcheckudata(L, 1, "_linearOutput", "torch.CudaTensor");
  THFloatTensor *buffer = luaT_getfieldcheckudata(L, 1, "_buffer", "torch.FloatTensor");
  
  THCudaTensor *
  THFloatTensor *
  
  THCudaTensor *logsoftOutput = luaT_getfieldcheckudata(L, 1, "_logSoftMaxOutput", "torch.CudaTensor");
  
  THCudaTensor *weight = luaT_getfieldcheckudata(L, 1, "weight", "torch.CudaTensor");
