     }
     childId = parentId;
   }
-  
-  while (1) 
-  {
-    float z = input_k[i];
-    float *nodeWeight = weight+
-    if(buffer[threadIdx.x] < z)
-      buffer[threadIdx.x] = z;
-  }
-
+  *output_k = narrowsum;
 }
 
 static int cunnx_SoftMaxTree_updateOutput(lua_State *L)

  int i_step = blockDim.x;
  int k = blockIdx.x;
  float *input_k = input + k*nInput;
  float *output_k = output + k*nInput;
  int childId = (*(target+k)) - 1;
  int parentId, parentIdx, childIdx, nChildren;
  int nOutput;

      childId = parentId;
    }
  }
  THIntTensor_free(node);
  
  // we use these to keep intermediate results for later backprop
  THCudaTensor_resize2d(linearOutput, n);
  THCudaTensor_resize1d(logsoftOutput, n);
  
  
  /* call cudakernel */
  dim3 blocks(input->size[0]); // each block is an example
  dim3 threads(SOFTMAXTREE_THREADS);
  cunnx_SoftMaxTree_updateOutput_kernel<<<blocks,threads>>>(
    THCudaTensor_data(output), THCudaTensor_data(linearOutput), 
    THCudaTensor_data(logsoftOutput), THCudaTensor_data(input), 
    THCudaTensor_data(weight), THCudaTensor_data(bias), 
    THCudaTensor_data(target), THCudaTensor_data(childParent), 
    THCudaTensor_data(parentChildren), input->size[1], root_id
  );
  
  cudaError errcode = cudaGetLastError();
  if(errcode != cudaSuccess)
