 
 
 __global__ void cunnx_SoftMaxTree_updateGradInput_kernel(
-  float *gradInput, float* logsoftOutput, float *gradOutput, float* weight,
-  float *target, float* childParent, float* parentChildren, 
+  float *gradInput, float *logsoftOutput, float *gradOutput, float* weight,
+  float *target, float *childParent, float *parentChildren, 
   int nInput, int rootId, int maxFamilyPath)
 {
   //__shared__ float input_buffer[nInput]; // constant might be faster



__global__ void cunnx_SoftMaxTree_updateOutput_kernel(
  float *output, float* logsoftOutput,
  float *input, float* weight, float* bias, 
  float *target, float* childParent, float* parentChildren, 
  int nInput, int rootId, int maxFamilyPath)
{
  //__shared__ float input_buffer[nInput]; // constant might be faster

  cunnx_SoftMaxTree_accGradParameters_kernel<<<blocks,threads>>>(
    THCudaTensor_data(gradWeight), THCudaTensor_data(Bias), 
    THCudaTensor_data(input), THCudaTensor_data(linearGradOutput), 
    THCudaTensor_data(nodeUpdateCuda), THCudaTensor_data(target), 
    THCudaTensor_data(childParent), THCudaTensor_data(parentChildren), 
    input->size[1], rootId, maxFamilyPath, maxDept, scale 
  );
  
  // copy nodeIds from device to host
  
  /* updates will contain parentId (key) sum of scales (value)*/
  lua_pushinteger(L, (int)(parentId+1));
  lua_gettable(L, -2);
  double count = lua_tonumber(L, -1) + scale;
  lua_pop(L, 1);
  
  lua_pushinteger(L, (int)(parentId+1)); /* key */
  lua_pushnumber(L, count); /* value */
  lua_settable(L, -3);
    
  return 0;
}
